[OPENMP50]Allow overlapping mapping in target constrcuts.

[ABataev]

Summary

OpenMP 5.0 removed a lot of restriction for overlapped mapped items
comparing to OpenMP 4.5. Patch restricts the checks for overlapped data
mappings only for OpenMP 4.5 and less and reorders mapping of the
arguments so, that present and alloc mappings are processed first and
then all others.

[ABataev]

Rebase + fixed implicit mapping of partially mapped structs.

[ABataev]

Rebase

[ABataev]

1. Generalized interface for all mapping operations.
2. Fixed compatibility of declare mapper and other mapping operations with OpenMP 5.0 examples.
3. Simplified declare mapper codegen.

[ABataev]

Rebase

[ABataev]

Rebase

[ABataev]

Rebase.
This patch is an important part of 'default' mappers lookup for data members.

[jdenny]

Thanks for working on this. Sorry to take so long to review. Before I try to digest the code, I have a few high-level questions.

Based on the test suite changes, TARGET_PARAM is disappearing from many cases. Can you explain a bit how that supports overlapping and reordering?

Have you considered issue 2337 for the OpenMP spec and how your implementation handles the ambiguous cases cited there?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7344

Unrelated change?

[ABataev]

In D86119#2330310, @jdenny wrote:

Thanks for working on this. Sorry to take so long to review. Before I try to digest the code, I have a few high-level questions.

Based on the test suite changes, TARGET_PARAM is disappearing from many cases. Can you explain a bit how that supports overlapping and reordering?

TARGET_PARAM is used only for marking the data that should be passed to the kernel function as an argument. We just generate it in many cases but the runtime actually does not use them. Thу patch relies on this thing, otherwise, we may pass an incorrect number of arguments to the kernel functions.

Have you considered issue 2337 for the OpenMP spec and how your implementation handles the ambiguous cases cited there?

Can you provide the details about this issue?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7344

Well, in general, yes, I just tried to simplify the code, can commit it as a separate NFC.

[jdenny]

In D86119#2330339, @ABataev wrote:

In D86119#2330310, @jdenny wrote:

Thanks for working on this. Sorry to take so long to review. Before I try to digest the code, I have a few high-level questions.

Based on the test suite changes, TARGET_PARAM is disappearing from many cases. Can you explain a bit how that supports overlapping and reordering?

TARGET_PARAM is used only for marking the data that should be passed to the kernel function as an argument. We just generate it in many cases but the runtime actually does not use them. Thу patch relies on this thing, otherwise, we may pass an incorrect number of arguments to the kernel functions.

Is it reasonable to extract that change into a parent patch? That would at least make the test suite changes easier to follow.

Have you considered issue 2337 for the OpenMP spec and how your implementation handles the ambiguous cases cited there?

Can you provide the details about this issue?

It discusses ambiguities in the way the spec describes map clause ordering. Here's one example relevant to present:

#pragma omp target exit data map(from: x) map(present, release: x)

One statement in the spec says a map clause with from is effectively ordered before one with release. Another statement says a map clause with present is effectively ordered before one without. Which applies in the above example? In some discussions, people agreed the statement about present was intended to have precedence, but the spec doesn't say that. That resolution probably makes sense at entry to a region, but does it make sense at exit? Would it suppress from in this example? (Actually, should there be two reference counter decrements in this example or just one?)

These ambiguities are the reason I punted on this issue when implementing present. If we can come up with a reasonable implementation for all cases, perhaps we can propose a new wording for the spec.

[ABataev]

In D86119#2330601, @jdenny wrote:

In D86119#2330339, @ABataev wrote:

In D86119#2330310, @jdenny wrote:

Thanks for working on this. Sorry to take so long to review. Before I try to digest the code, I have a few high-level questions.

Based on the test suite changes, TARGET_PARAM is disappearing from many cases. Can you explain a bit how that supports overlapping and reordering?

TARGET_PARAM is used only for marking the data that should be passed to the kernel function as an argument. We just generate it in many cases but the runtime actually does not use them. Thу patch relies on this thing, otherwise, we may pass an incorrect number of arguments to the kernel functions.

Is it reasonable to extract that change into a parent patch? That would at least make the test suite changes easier to follow.

I'll see what I can do about it.

Have you considered issue 2337 for the OpenMP spec and how your implementation handles the ambiguous cases cited there?

Can you provide the details about this issue?

It discusses ambiguities in the way the spec describes map clause ordering. Here's one example relevant to present:

#pragma omp target exit data map(from: x) map(present, release: x)

One statement in the spec says a map clause with from is effectively ordered before one with release. Another statement says a map clause with present is effectively ordered before one without. Which applies in the above example? In some discussions, people agreed the statement about present was intended to have precedence, but the spec doesn't say that. That resolution probably makes sense at entry to a region, but does it make sense at exit? Would it suppress from in this example? (Actually, should there be two reference counter decrements in this example or just one?)

These ambiguities are the reason I punted on this issue when implementing present. If we can come up with a reasonable implementation for all cases, perhaps we can propose a new wording for the spec.

In tgis implementation, the mapping withthe present modifier will be ordered to be the first.

[ABataev]

Rebase

[ABataev]

Rebase

[grokos]

openmp/libomptarget/src/omptarget.cpp

238

What is the current status of the order of the arguments clang emits? Is it still necessary to traverse arguments in reverse order here?

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

Yes, still required

[JonChesterfield]

If data mappings can overlap, then it follows that copies to/from the target must be done sequentially by the runtime, unless additional information on their independence exists. Alias analysis style.

I see the programmer convenience angle, but it is a shame to no longer be able to perform the to: mappings simultaneously.

[ye-luo]

1. Could you separate the reordering related changes to separate patch?
2. Could you mention which line in spec 4.5 was the restriction? Even 5.0/5.1 has some restrictions. Need to be clear which one you refer to.

[ABataev]

In D86119#2402661, @ye-luo wrote:

1. Could you separate the reordering related changes to separate patch?

It is impossible. There is a separate patch for removing extra TGT_TARGET_PARAM but it cannot be separated from this patch.

2. Could you mention which line in spec 4.5 was the restriction? Even 5.0/5.1 has some restrictions. Need to be clear which one you refer to.

Yes, you right, I'm trying to support something like map(alloc: S) map(tofrom: S.f) or map(S, S.ptr[0]) which was not previously allowed and is a feature of OpenMP 5.0.

[ABataev]

Rebase

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

Based on the conversation in
https://reviews.llvm.org/D85216
This line of code neither before nor after the change plays well.
Shall we fix the order in targetDataEnd first?

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

This change is part of this patch and cannot be committed separately.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

I mean could you fix that issue as a parent of this patch?

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

This change is part of this patch and cannot be committed separately.

If fixing the reordering is part of this patch, I should have seen "target_data_function == targetDataEnd ?" branches disappear.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

Nope, just with this patch. It reorders the maps and need to change the cleanup order too.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

It works just like constructors/destructors: allocate in direct order, deallocate in reversed to correctly handle map order.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

The description says that "present and alloc mappings are processed first and then all others."
Why the order of arguments in targetDataBegin, targetDataEnd and targetDataUpdate all get reversed.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

Because this is for mappers. Mapper maps are ordered by the compiler in the direct order (alloc, maps, delete) but when we need to do exit, we need to release the data in reversed order (deletes, maps, allocs).

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

I was not making the question clear. My question about "reverse" is not about having a reverse order for targetDataBegin. My question was about "reversing" from the the old code. Your change put the opposite order for targetDataBegin, targetDataEnd and targetDataUpdate cases.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

I was not making the question clear. My question about "reverse" is not about having a reverse order for targetDataBegin. My question was about "reversing" from the the old code. Your change put the opposite order for targetDataBegin, targetDataEnd and targetDataUpdate cases.

typo correction
I was not making the question clear. My question about "reverse" is not about having a reverse order for targetDataEnd. My question was about "reversing" from the the old code. Your change put the opposite order for targetDataBegin, targetDataEnd and targetDataUpdate cases.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

My separate question specifically for targetDataEnd is the following.

In target(), we call

targetDataBegin(args)
{  // forward order
  for (int32_t i = 0; i < arg_num; ++i) { ... }
}
launch_kernels
targetDataEnd(args)
{  // reverse order
  for (int32_t I = ArgNum - 1; I >= 0; --I) { }
}

At a mapper,

targetDataMapper
{
  // generate args_reverse in reverse order for targetDataEnd
  targetDataEnd(args_reverse)
}

Are we actually getting the original forward order due to one reverse in targetDataMapper and second reverse in targetDataEnd? Is this the desired behavior? This part confused me. Do I miss something? Could you explain a bit?

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

Yes, something like this. targetDataEnd reverses the order of mapping arrays. But mapper generator always generates mapping arrays in the direct order (it fills mapping arrays that later processed by the targetDataEnd function). We could fix this by passing extra Boolean flag to the generator function but it means the redesign of the mappers. That's why we have to reverse it in the libomptarget.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

You can check it yourself. Apply the patch, restore the original behavior in libomptarget and run libomptarget tests. Mapper related tests will crash.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

Stick with mapper generator always generating mapping arrays in the direct order. The targetDataMapper reverse the mapping array and then passes args_reverse into targetDataEnd. Inside targetDataEnd, mapping

238

Yes, something like this. targetDataEnd reverses the order of mapping arrays. But mapper generator always generates mapping arrays in the direct order (it fills mapping arrays that later processed by the targetDataEnd function). We could fix this by passing extra Boolean flag to the generator function but it means the redesign of the mappers. That's why we have to reverse it in the libomptarget.

Stick with mapper generator always generating mapping arrays in the direct order.

In the targetDataBegin case, targetDataMapper keep direct order args and calls targetDataBegin(args) and targetDataBegin process args in direct order.

In the targetDataEnd case, targetDataMapper reverses the mapping array and then passes args_reverse into targetDataEnd. Inside targetDataEnd, args_reverse are processed in reverse order. So targetDataEnd is actually processing the args in original direct order. This seems contradictory to the constructor/deconstructor like behavior that all the mappings must be processed in the actual reverse order in targetDataEnd.

This is my understanding. The current code should be wrong but obviously the current code is working. So why the current code is working? what is inconsistent in my analysis. Could you point out the missing piece.

238

You can check it yourself. Apply the patch, restore the original behavior in libomptarget and run libomptarget tests. Mapper related tests will crash.

For sure without this line, tests would crash and that is why you included this line of change in the patch. Since you made the change, you could explain why, right?

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

I changed and simplified codegen for the mapper generator without changing its interface. I could do this because of the new ordering, before we could not rely on it. But it also requires a change in the runtime.
targetDataEnd calls targetDataMapper and targetDataMapper fills the array in the direct order, but targetDataEnd processes them in the reverse order, but mapper generator does not know about it. It also has to generate the data in the reverse order, just like targetDataEnd does.

Before this patch mapper generator tried to do some ordering but it was not always correct. It was not expecting something like map(alloc:s) map(s.a) because it was not allowed by the compiler. That's why it worked before and won't work with this patch.
PS. The change in the mapper generator is also required and cannot be separated. Without this mappers tests won't work.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

I played a bit with your patch.

#pragma omp target exit data map(from: c.a[0:NUM], c.b[0:NU2M]) map(delete: c)

I put NUM=1024 and NU2M = 2048.
LIBOMPTARGET_DEBUG reports

Libomptarget --> Entry  0: Base=0x00007fff064080a8, Begin=0x00007fff064080a8, Size=16, Type=0x0, Name=(null)
Libomptarget --> Entry  1: Base=0x00007fff064080a8, Begin=0x0000000000f9cbd0, Size=4096, Type=0x1000000000012, Name=(null)
Libomptarget --> Entry  2: Base=0x00007fff064080b0, Begin=0x0000000000f86e10, Size=8192, Type=0x1000000000012, Name=(null)
Libomptarget --> Entry  3: Base=0x00007fff064080a8, Begin=0x00007fff064080a8, Size=16, Type=0x1000000000008, Name=(null)

Since targetDataEnd internally reverse the processing order, could you confirm that the frontend was emitting entries 3,2,1,0?
I'm wondering if the frontend could emit 3, 0, 1, 2 so the processing order is 2,1,0,3? The spec requires struct element processed before the struct in "target exit data"

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

No, the frontend emits in the order 0, 1, 2, 3. targetDataEnd process in reversed order 3, 2, 1, 0, but the mapper does not know about it and still emits the data in the order 0, 1, 2, 3.
And it is only for mappers!
So, say you have an extra map something like map(a).

map (a) map(mapper(id), tofrom: c)

where mapper for с does something like you wrote.

In this case the order would be 0, 1, 2, 3, 4, where 0 is mapping of a and 1-4 is mapping of c.
When we need to delete the data, the mapper still would generate 1,2,3,4 + 0 for mapping of a, but targetDataEnd expects 4,3,2,1,0. That's why we have to reverse the mapping data, produced by the mapper generator for targetDataEnd.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

Let us leave the mapper case aside which has extra mess.
Double checked that "Libomptarget --> Entry 0" is printed at the __tgt_target_data_end_mapper. So the order is as you said 0, 1, 2, 3 from the frontend. What is the "Entry 0"? more specifically is the difference between entry 0 and 3? Entry 0 seems to be an implicit map while 3 is explicit.

#pragma omp target exit data map(from: c.a[0:NUM], c.b[0:NU2M]) map(delete: c)
the "map(delete: c)" has some state machine to protect the delete due to ordering.
So I'm wondering why the frontend must issue both 0 and 3. Can the front end fuse 0 and 3?
I mean the frontend generates 3, 1, 2 and the runtime processing 2,1,3 without the deleting issue?

[ABataev]

This comment has been deleted.

openmp/libomptarget/src/omptarget.cpp

238

Entry 0 is the address that should be passed to the kernel (for the captured variable, that's why it is marked as TGT_TARGET_PARAM - target kernel argument), entries 1-3 are the actual mappings. Yes, I assume the frontend can fuse these entries in many cases, but it is different problem that should be addressed in a separate patch.

[ye-luo]

openmp/libomptarget/src/omptarget.cpp

238

There is no kernel associated enter/exit data...

Thanks for answering all my puzzles. If you think some of my questions can be better answered by some documentation, please point me if there are any.

I think both the frontend and runtime needs to be further consolidated to reduce side effects in future patches. The current patch in the runtime library part looks good to me for the current need.

Please ping appropriate reviewers for the frontend change, so we can keep this patch moving.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

No problem!
Yes, there is no kernel, missed it, but still it is the reason. The codegen you see is not directly related to this patch, this is just how the mapping currently works, it is just not quite optimal. Sure, it can be improved in many cases but just like I said it is different problem that should be addressed in separate patch(es). Also, I believe some of the optimizations can be implemented in OpenMPOpt pass.

[ABataev]

openmp/libomptarget/src/omptarget.cpp

238

For the documentations, there is design description for the mappers https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx. For the mapping rules, see the comments in generateInfoForComponentList function

[cchen]

Just tried this patch and found that below test abort inside generateInfoCapture.

#include <stdio.h>
#define LEN 100
int buf[LEN];

int main()
{
    int i;
    int *p = buf;

    for (i = 0; i < LEN; i++) {
        p[i] = 0;
    }

#pragma omp target map(p) map(p[:100])
    {
        *(p+5) = 1;
    }

}

[ABataev]

In D86119#2437676, @cchen wrote:

Just tried this patch and found that below test abort inside generateInfoCapture.

#include <stdio.h>
#define LEN 100
int buf[LEN];

int main()
{
    int i;
    int *p = buf;

    for (i = 0; i < LEN; i++) {
        p[i] = 0;
    }

#pragma omp target map(p) map(p[:100])
    {
        *(p+5) = 1;
    }

}

I'll check this, thanks!

[ABataev]

Fix analysis for overlapping pointers.

[cchen]

Below test asserts inside generateInfoForComponentList:

Assertion failed: (!IsPointer && "Unexpected base element with the pointer type."), function generateInfoForComponentList, file /Users/cchen/workspace/llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp, line 7754.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.

#define LEN 100

struct Type {
    int *p;
};

int buf[LEN];

int main()
{
    int i;
    struct Type s;
    s.p = buf;

    for (i = 0; i < LEN; i++) {
        s.p[i] = 0;
    }

#pragma omp target map(s.p) map(s.p[:100])
    {
        *(s.p+5) = 1;
    }
}

[ABataev]

General approach to fix the analysis of pointers.

[ABataev]

Update

[ABataev]

Update+rebase, need to update the tests.

[ABataev]

Fixed condition logic.

[ABataev]

Updated tests

[abhinavgaba]

Thanks for the changes, Alexey! I tried the patch locally, and it looks stable. It handled several tests I tried, including the following case involving array section on a pointer to pointer base, and nested mappers with PTR_AND_OBJ maps successfully:

#include <stdio.h>

typedef struct { int a; double *b; } C;
#pragma omp declare mapper(id1: C s) map(to:s.a) map(from:s.b[0:2])

typedef struct { int e; C f; int h; short *g; } D;
#pragma omp declare mapper(default: D r) map(from:r.e) map(mapper(id1), tofrom:r.f) map(tofrom: r.g[0:r.h])

int main() {
  constexpr int N = 10;
  D s;
  s.e = 111;
  s.f.a = 222;
  double x[2]; x[1] = 20;
  short y[N]; y[1] = 30;
  s.f.b = &x[0];
  s.g = &y[0];
  s.h = N;

  D* sp = &s;
  D** spp = &sp;

  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 111 222 20.0 <host_addr1> 30 <host_addr2>

  #pragma omp target map(tofrom:spp[0][0])
  {
    printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
    // Expected: <not 111> 222 <not 20.0> <dev_addr1> 30 <dev_addr2>
    spp[0][0].e = 333;
    spp[0][0].f.a = 444;
    spp[0][0].f.b[1] = 40;
    spp[0][0].g[1] = 50;
  }
  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 333 222 40.0 <host_addr1> 50 <host_addr2>
}

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8483–8484

Tabs should probably be spaces. Same for a few other places in the changeset.

9566–9567

IsDelete

9729

Commented-out code intentionally left in?

[grokos]

In D86119#2561163, @abhinavgaba wrote:

Thanks for the changes, Alexey! I tried the patch locally, and it looks stable. It handled several tests I tried, including the following case involving array section on a pointer to pointer base, and nested mappers with PTR_AND_OBJ maps successfully:

#include <stdio.h>

typedef struct { int a; double *b; } C;
#pragma omp declare mapper(id1: C s) map(to:s.a) map(from:s.b[0:2])

typedef struct { int e; C f; int h; short *g; } D;
#pragma omp declare mapper(default: D r) map(from:r.e) map(mapper(id1), tofrom:r.f) map(tofrom: r.g[0:r.h])

int main() {
  constexpr int N = 10;
  D s;
  s.e = 111;
  s.f.a = 222;
  double x[2]; x[1] = 20;
  short y[N]; y[1] = 30;
  s.f.b = &x[0];
  s.g = &y[0];
  s.h = N;

  D* sp = &s;
  D** spp = &sp;

  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 111 222 20.0 <host_addr1> 30 <host_addr2>

  #pragma omp target map(tofrom:spp[0][0])
  {
    printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
    // Expected: <not 111> 222 <not 20.0> <dev_addr1> 30 <dev_addr2>
    spp[0][0].e = 333;
    spp[0][0].f.a = 444;
    spp[0][0].f.b[1] = 40;
    spp[0][0].g[1] = 50;
  }
  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 333 222 40.0 <host_addr1> 50 <host_addr2>
}

@ABataev This is a nice complex example, I think it's worth including it in the runtime tests (under libomptarget).

@abhinavgaba Thanks for providing it!

[ABataev]

In D86119#2562187, @grokos wrote:

In D86119#2561163, @abhinavgaba wrote:

Thanks for the changes, Alexey! I tried the patch locally, and it looks stable. It handled several tests I tried, including the following case involving array section on a pointer to pointer base, and nested mappers with PTR_AND_OBJ maps successfully:

#include <stdio.h>

typedef struct { int a; double *b; } C;
#pragma omp declare mapper(id1: C s) map(to:s.a) map(from:s.b[0:2])

typedef struct { int e; C f; int h; short *g; } D;
#pragma omp declare mapper(default: D r) map(from:r.e) map(mapper(id1), tofrom:r.f) map(tofrom: r.g[0:r.h])

int main() {
  constexpr int N = 10;
  D s;
  s.e = 111;
  s.f.a = 222;
  double x[2]; x[1] = 20;
  short y[N]; y[1] = 30;
  s.f.b = &x[0];
  s.g = &y[0];
  s.h = N;

  D* sp = &s;
  D** spp = &sp;

  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 111 222 20.0 <host_addr1> 30 <host_addr2>

  #pragma omp target map(tofrom:spp[0][0])
  {
    printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
    // Expected: <not 111> 222 <not 20.0> <dev_addr1> 30 <dev_addr2>
    spp[0][0].e = 333;
    spp[0][0].f.a = 444;
    spp[0][0].f.b[1] = 40;
    spp[0][0].g[1] = 50;
  }
  printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
  // Expected: 333 222 40.0 <host_addr1> 50 <host_addr2>
}

@ABataev This is a nice complex example, I think it's worth including it in the runtime tests (under libomptarget).

@abhinavgaba Thanks for providing it!

Ok, will add it as a part of the patch

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8483–8484

These are not tabs. Looks like this is how Phabricators shows some format changes.

9729

Yeah, forgot to remove it, thanks.

[ABataev]

Fixes, added a test

[abhinavgaba]

Thanks for the changes, Alexey.

clang/lib/CodeGen/CGOpenMPRuntime.h

void emitThreadPrivateVarInit(CodeGenFunction &CGF, Address VDAddr,
                                 llvm::Value *Dtor, SourceLocation Loc);
 
   /// Emit the array initialization or deletion portion for user-defined mapper
   /// code generation.
   void emitUDMapperArrayInitOrDel(CodeGenFunction &MapperCGF,
                                   llvm::Value *Handle, llvm::Value *BasePtr,
                                   llvm::Value *Ptr, llvm::Value *Size,
                                   llvm::Value *MapType, CharUnits ElementSize,
-                                  llvm::BasicBlock *ExitBB, bool IsInit);
+                                  llvm::BasicBlock *ExitBB);
 
   struct TaskResultTy {
     llvm::Value *NewTask = nullptr;
     llvm::Function *TaskEntry = nullptr;
     llvm::Value *NewTaskNewTaskTTy = nullptr;
     LValue TDBase;
     const RecordDecl *KmpTaskTQTyRD = nullptr;
     llvm::Value *TaskDupFn = nullptr;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

public:
   ///                    HE(FieldIndex, Pointer)}
   struct StructRangeInfoTy {
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> LowestElem = {
         0, Address::invalid()};
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> HighestElem = {
         0, Address::invalid()};
     Address Base = Address::invalid();
     bool IsArraySection = false;
+    bool HasCompleteRecord = false;
   };
 
 private:
   /// Kind that defines how a device pointer has to be returned.
   struct MapInfo {
     OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
     ArrayRef<OpenMPMapModifierKind> MapModifiers;
OpenMPOffloadMappingFlags getMapTypeBits(
     if (AddIsTargetParamFlag)
       Bits |= OMP_MAP_TARGET_PARAM;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_always)
         != MapModifiers.end())
       Bits |= OMP_MAP_ALWAYS;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
         != MapModifiers.end())
       Bits |= OMP_MAP_CLOSE;
-    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
-        != MapModifiers.end())
-      Bits |= OMP_MAP_PRESENT;
-    if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present)
-        != MotionModifiers.end())
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present) !=
+            MapModifiers.end() ||
+        llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) !=
+            MotionModifiers.end())
       Bits |= OMP_MAP_PRESENT;
     if (IsNonContiguous)
       Bits |= OMP_MAP_NON_CONTIG;
     return Bits;
   }
 
   /// Return true if the provided expression is a final array section. A
   /// final array section, is one whose length can't be proved to be one.
for (; I != CE; ++I) {
         QualType Ty = Components.rbegin()
                           ->getAssociatedDeclaration()
                           ->getType()
                           .getNonReferenceType();
         BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
         FirstPointerInComplexData = false;
       }
     }
+    // If ran into the whole component - allocate the space for the whole
+    // record.
+    if (!EncounteredME)
+      PartialStruct.HasCompleteRecord = true;
 
     if (!IsNonContiguous)
       return;
 
     const ASTContext &Context = CGF.getContext();
 
     // For supporting stride in array section, we need to initialize the first
     // dimension size as 1, first offset as 0, and first count as 1
for (const llvm::PointerUnion<const CXXRecordDecl *, const FieldDecl *>
         continue;
       if (const auto *Base = Data.dyn_cast<const CXXRecordDecl *>())
         getPlainLayout(Base, Layout, /*AsBase=*/true);
       else
         Layout.push_back(Data.get<const FieldDecl *>());
     }
   }
 
-public:
-  MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
-      : CurDir(&Dir), CGF(CGF) {
-    // Extract firstprivate clause information.
-    for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
-      for (const auto *D : C->varlists())
-        FirstPrivateDecls.try_emplace(
-            cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl()), C->isImplicit());
-    // Extract implicit firstprivates from uses_allocators clauses.
-    for (const auto *C : Dir.getClausesOfKind<OMPUsesAllocatorsClause>()) {
-      for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
-        OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
-        if (const auto *DRE = dyn_cast_or_null<DeclRefExpr>(D.AllocatorTraits))
-          FirstPrivateDecls.try_emplace(cast<VarDecl>(DRE->getDecl()),
-                                        /*Implicit=*/true);
-        else if (const auto *VD = dyn_cast<VarDecl>(
-                     cast<DeclRefExpr>(D.Allocator->IgnoreParenImpCasts())
-                         ->getDecl()))
-          FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
-      }
-    }
-    // Extract device pointer clause information.
-    for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
-      for (auto L : C->component_lists())
-        DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
-  }
-
-  /// Constructor for the declare mapper directive.
-  MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
-      : CurDir(&Dir), CGF(CGF) {}
-
-  /// Generate code for the combined entry if we have a partially mapped struct
-  /// and take care of the mapping flags of the arguments corresponding to
-  /// individual struct members.
-  void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
-                         MapFlagsArrayTy &CurTypes,
-                         const StructRangeInfoTy &PartialStruct,
-                         const ValueDecl *VD = nullptr,
-                         bool NotTargetParams = true) const {
-    if (CurTypes.size() == 1 &&
-        ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) &&
-        !PartialStruct.IsArraySection)
-      return;
-    CombinedInfo.Exprs.push_back(VD);
-    // Base is the base of the struct
-    CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer());
-    // Pointer is the address of the lowest element
-    llvm::Value *LB = PartialStruct.LowestElem.second.getPointer();
-    CombinedInfo.Pointers.push_back(LB);
-    // There should not be a mapper for a combined entry.
-    CombinedInfo.Mappers.push_back(nullptr);
-    // Size is (addr of {highest+1} element) - (addr of lowest element)
-    llvm::Value *HB = PartialStruct.HighestElem.second.getPointer();
-    llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1);
-    llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy);
-    llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy);
-    llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr);
-    llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty,
-                                                  /*isSigned=*/false);
-    CombinedInfo.Sizes.push_back(Size);
-    // Map type is always TARGET_PARAM, if generate info for captures.
-    CombinedInfo.Types.push_back(NotTargetParams ? OMP_MAP_NONE
-                                                 : OMP_MAP_TARGET_PARAM);
-    // If any element has the present modifier, then make sure the runtime
-    // doesn't attempt to allocate the struct.
-    if (CurTypes.end() !=
-        llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) {
-          return Type & OMP_MAP_PRESENT;
-        }))
-      CombinedInfo.Types.back() |= OMP_MAP_PRESENT;
-    // Remove TARGET_PARAM flag from the first element if any.
-    if (!CurTypes.empty())
-      CurTypes.front() &= ~OMP_MAP_TARGET_PARAM;
-
-    // All other current entries will be MEMBER_OF the combined entry
-    // (except for PTR_AND_OBJ entries which do not have a placeholder value
-    // 0xFFFF in the MEMBER_OF field).
-    OpenMPOffloadMappingFlags MemberOfFlag =
-        getMemberOfFlag(CombinedInfo.BasePointers.size() - 1);
-    for (auto &M : CurTypes)
-      setCorrectMemberOfFlag(M, MemberOfFlag);
-  }
-
   /// Generate all the base pointers, section pointers, sizes, map types, and
   /// mappers for the extracted mappable expressions (all included in \a
   /// CombinedInfo). Also, for each item that relates with a device pointer, a
   /// pair of the relevant declaration and index where it occurs is appended to
   /// the device pointers info array.
-  void generateAllInfo(
-      MapCombinedInfoTy &CombinedInfo,
+  void generateAllInfoForClauses(
+      ArrayRef<const OMPClause *> Clauses, MapCombinedInfoTy &CombinedInfo,
       const llvm::DenseSet<CanonicalDeclPtr<const Decl>> &SkipVarSet =
           llvm::DenseSet<CanonicalDeclPtr<const Decl>>()) const {
     // We have to process the component lists that relate with the same
     // declaration in a single chunk so that we can generate the map flags
     // correctly. Therefore, we organize all lists in a map.
-    llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
+    enum MapKind { Present, Allocs, Other, Total };
+    llvm::MapVector<CanonicalDeclPtr<const Decl>,
+                    SmallVector<SmallVector<MapInfo, 8>, 4>>
+        Info;
 
     // Helper function to fill the information map for the different supported
     // clauses.
     auto &&InfoGen =
         [&Info, &SkipVarSet](
-            const ValueDecl *D,
+            const ValueDecl *D, MapKind Kind,
             OMPClauseMappableExprCommon::MappableExprComponentListRef L,
             OpenMPMapClauseKind MapType,
             ArrayRef<OpenMPMapModifierKind> MapModifiers,
             ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
             bool ReturnDevicePointer, bool IsImplicit, const ValueDecl *Mapper,
             const Expr *VarRef = nullptr, bool ForDeviceAddr = false) {
-          const ValueDecl *VD =
-              D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
-          if (SkipVarSet.count(VD))
-            return;
-          Info[VD].emplace_back(L, MapType, MapModifiers, MotionModifiers,
-                                ReturnDevicePointer, IsImplicit, Mapper, VarRef,
-                                ForDeviceAddr);
-        };
-
-    assert(CurDir.is<const OMPExecutableDirective *>() &&
-           "Expect a executable directive");
-    const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
-    for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
+          if (SkipVarSet.contains(D))
+            return;
+          auto It = Info.find(D);
+          if (It == Info.end())
+            It = Info
+                     .insert(std::make_pair(
+                         D, SmallVector<SmallVector<MapInfo, 8>, 4>(Total)))
+                     .first;
+          It->second[Kind].emplace_back(
+              L, MapType, MapModifiers, MotionModifiers, ReturnDevicePointer,
+              IsImplicit, Mapper, VarRef, ForDeviceAddr);
+        };
+
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPMapClause>(Cl);
+      if (!C)
+        continue;
+      MapKind Kind = Other;
+      if (!C->getMapTypeModifiers().empty() &&
+          llvm::any_of(C->getMapTypeModifiers(), [](OpenMPMapModifierKind K) {
+            return K == OMPC_MAP_MODIFIER_present;
+          }))
+        Kind = Present;
+      else if (C->getMapType() == OMPC_MAP_alloc)
+        Kind = Allocs;
       const auto *EI = C->getVarRefs().begin();
       for (const auto L : C->component_lists()) {
-        // The Expression is not correct if the mapping is implicit
         const Expr *E = (C->getMapLoc().isValid()) ? *EI : nullptr;
-        InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(),
+        InfoGen(std::get<0>(L), Kind, std::get<1>(L), C->getMapType(),
                 C->getMapTypeModifiers(), llvm::None,
                 /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L),
                 E);
         ++EI;
       }
     }
-    for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>()) {
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPToClause>(Cl);
+      if (!C)
+        continue;
+      MapKind Kind = Other;
+      if (!C->getMotionModifiers().empty() &&
+          llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) {
+            return K == OMPC_MOTION_MODIFIER_present;
+          }))
+        Kind = Present;
       const auto *EI = C->getVarRefs().begin();
       for (const auto L : C->component_lists()) {
-        InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None,
+        InfoGen(std::get<0>(L), Kind, std::get<1>(L), OMPC_MAP_to, llvm::None,
                 C->getMotionModifiers(), /*ReturnDevicePointer=*/false,
                 C->isImplicit(), std::get<2>(L), *EI);
         ++EI;
       }
     }
-    for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>()) {
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPFromClause>(Cl);
+      if (!C)
+        continue;
+      MapKind Kind = Other;
+      if (!C->getMotionModifiers().empty() &&
+          llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) {
+            return K == OMPC_MOTION_MODIFIER_present;
+          }))
+        Kind = Present;
       const auto *EI = C->getVarRefs().begin();
       for (const auto L : C->component_lists()) {
-        InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None,
+        InfoGen(std::get<0>(L), Kind, std::get<1>(L), OMPC_MAP_from, llvm::None,
                 C->getMotionModifiers(), /*ReturnDevicePointer=*/false,
                 C->isImplicit(), std::get<2>(L), *EI);
         ++EI;
       }
     }
 
     // Look at the use_device_ptr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_ptr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
     // emission of that entry until the whole struct has been processed.
-    llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
+    llvm::MapVector<CanonicalDeclPtr<const Decl>,
+                    SmallVector<DeferredDevicePtrEntryTy, 4>>
         DeferredInfo;
     MapCombinedInfoTy UseDevicePtrCombinedInfo;
 
-    for (const auto *C :
-         CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
+      if (!C)
+        continue;
       for (const auto L : C->component_lists()) {
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
             std::get<1>(L);
         assert(!Components.empty() &&
                "Not expecting empty list of components!");
         const ValueDecl *VD = Components.back().getAssociatedDeclaration();
         VD = cast<ValueDecl>(VD->getCanonicalDecl());
         const Expr *IE = Components.back().getAssociatedExpression();
         // If the first component is a member expression, we have to look into
         // 'this', which maps to null in the map of map information. Otherwise
         // look directly for the information.
         auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
 
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
-          auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
+          bool Found = false;
+          for (auto &Data : It->second) {
+            auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
               return MI.Components.back().getAssociatedDeclaration() == VD;
             });
-          // If we found a map entry, signal that the pointer has to be returned
-          // and move on to the next declaration.
-          // Exclude cases where the base pointer is mapped as array subscript,
-          // array section or array shaping. The base address is passed as a
-          // pointer to base in this case and cannot be used as a base for
-          // use_device_ptr list item.
-          if (CI != It->second.end()) {
+            // If we found a map entry, signal that the pointer has to be
+            // returned and move on to the next declaration. Exclude cases where
+            // the base pointer is mapped as array subscript, array section or
+            // array shaping. The base address is passed as a pointer to base in
+            // this case and cannot be used as a base for use_device_ptr list
+            // item.
+            if (CI != Data.end()) {
               auto PrevCI = std::next(CI->Components.rbegin());
               const auto *VarD = dyn_cast<VarDecl>(VD);
               if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
                   isa<MemberExpr>(IE) ||
                   !VD->getType().getNonReferenceType()->isPointerType() ||
                   PrevCI == CI->Components.rend() ||
                   isa<MemberExpr>(PrevCI->getAssociatedExpression()) || !VarD ||
                   VarD->hasLocalStorage()) {
                 CI->ReturnDevicePointer = true;
-              continue;
+                Found = true;
+                break;
+              }
             }
           }
+          if (Found)
+            continue;
         }
 
         // We didn't find any match in our map information - generate a zero
         // size array section - if the pointer is a struct member we defer this
         // action until the whole struct has been processed.
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
-          InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None, llvm::None,
-                  /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
+          InfoGen(nullptr, Other, Components, OMPC_MAP_unknown, llvm::None,
+                  llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
+                  nullptr);
           DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false);
         } else {
           llvm::Value *Ptr =
               CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
           UseDevicePtrCombinedInfo.Exprs.push_back(VD);
           UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD);
           UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
           UseDevicePtrCombinedInfo.Sizes.push_back(
               llvm::Constant::getNullValue(CGF.Int64Ty));
           UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
           UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
         }
       }
     }
 
     // Look at the use_device_addr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_addr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
     // emission of that entry until the whole struct has been processed.
     llvm::SmallDenseSet<CanonicalDeclPtr<const Decl>, 4> Processed;
-    for (const auto *C :
-         CurExecDir->getClausesOfKind<OMPUseDeviceAddrClause>()) {
+    for (const auto *Cl : Clauses) {
+      const auto *C = dyn_cast<OMPUseDeviceAddrClause>(Cl);
+      if (!C)
+        continue;
       for (const auto L : C->component_lists()) {
         assert(!std::get<1>(L).empty() &&
                "Not expecting empty list of components!");
         const ValueDecl *VD = std::get<1>(L).back().getAssociatedDeclaration();
         if (!Processed.insert(VD).second)
           continue;
         VD = cast<ValueDecl>(VD->getCanonicalDecl());
         const Expr *IE = std::get<1>(L).back().getAssociatedExpression();
         // If the first component is a member expression, we have to look into
         // 'this', which maps to null in the map of map information. Otherwise
         // look directly for the information.
         auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
 
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
-          auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
+          bool Found = false;
+          for (auto &Data : It->second) {
+            auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
               return MI.Components.back().getAssociatedDeclaration() == VD;
             });
-          // If we found a map entry, signal that the pointer has to be returned
-          // and move on to the next declaration.
-          if (CI != It->second.end()) {
+            // If we found a map entry, signal that the pointer has to be
+            // returned and move on to the next declaration.
+            if (CI != Data.end()) {
               CI->ReturnDevicePointer = true;
-            continue;
+              Found = true;
+              break;
             }
           }
+          if (Found)
+            continue;
+        }
 
         // We didn't find any match in our map information - generate a zero
         // size array section - if the pointer is a struct member we defer this
         // action until the whole struct has been processed.
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
-          InfoGen(nullptr, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
+          InfoGen(nullptr, Other, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
                   llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
                   nullptr, nullptr, /*ForDeviceAddr=*/true);
           DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true);
         } else {
           llvm::Value *Ptr;
           if (IE->isGLValue())
             Ptr = CGF.EmitLValue(IE).getPointer(CGF);
           else
             Ptr = CGF.EmitScalarExpr(IE);
           CombinedInfo.Exprs.push_back(VD);
           CombinedInfo.BasePointers.emplace_back(Ptr, VD);
           CombinedInfo.Pointers.push_back(Ptr);
           CombinedInfo.Sizes.push_back(
               llvm::Constant::getNullValue(CGF.Int64Ty));
           CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
           CombinedInfo.Mappers.push_back(nullptr);
         }
       }
     }
 
-    for (const auto &M : Info) {
-      // Underlying variable declaration used in the map clause.
-      const ValueDecl *VD = std::get<0>(M);
-
+    for (const auto &Data : Info) {
+      StructRangeInfoTy PartialStruct;
       // Temporary generated information.
       MapCombinedInfoTy CurInfo;
-      StructRangeInfoTy PartialStruct;
-
-      for (const MapInfo &L : M.second) {
+      const Decl *D = Data.first;
+      const ValueDecl *VD = cast_or_null<ValueDecl>(D);
+      for (const auto &M : Data.second) {
+        for (const MapInfo &L : M) {
           assert(!L.Components.empty() &&
                  "Not expecting declaration with no component lists.");
 
           // Remember the current base pointer index.
           unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size();
           CurInfo.NonContigInfo.IsNonContiguous =
               L.Components.back().isNonContiguous();
           generateInfoForComponentList(
-            L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo,
-            PartialStruct, /*IsFirstComponentList=*/false, L.IsImplicit,
-            L.Mapper, L.ForDeviceAddr, VD, L.VarRef);
+              L.MapType, L.MapModifiers, L.MotionModifiers, L.Components,
+              CurInfo, PartialStruct, /*IsFirstComponentList=*/false,
+              L.IsImplicit, L.Mapper, L.ForDeviceAddr, VD, L.VarRef);
 
           // If this entry relates with a device pointer, set the relevant
           // declaration and add the 'return pointer' flag.
           if (L.ReturnDevicePointer) {
             assert(CurInfo.BasePointers.size() > CurrentBasePointersIdx &&
                    "Unexpected number of mapped base pointers.");
 
             const ValueDecl *RelevantVD =
                 L.Components.back().getAssociatedDeclaration();
             assert(RelevantVD &&
                    "No relevant declaration related with device pointer??");
 
             CurInfo.BasePointers[CurrentBasePointersIdx].setDevicePtrDecl(
                 RelevantVD);
             CurInfo.Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PARAM;
           }
         }
+      }
 
       // Append any pending zero-length pointers which are struct members and
       // used with use_device_ptr or use_device_addr.
-      auto CI = DeferredInfo.find(M.first);
+      auto CI = DeferredInfo.find(Data.first);
       if (CI != DeferredInfo.end()) {
         for (const DeferredDevicePtrEntryTy &L : CI->second) {
           llvm::Value *BasePtr;
           llvm::Value *Ptr;
           if (L.ForDeviceAddr) {
             if (L.IE->isGLValue())
               Ptr = this->CGF.EmitLValue(L.IE).getPointer(CGF);
             else
               Ptr = this->CGF.EmitScalarExpr(L.IE);
             BasePtr = Ptr;
             // Entry is RETURN_PARAM. Also, set the placeholder value
             // MEMBER_OF=FFFF so that the entry is later updated with the
             // correct value of MEMBER_OF.
             CurInfo.Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_MEMBER_OF);
           } else {
             BasePtr = this->CGF.EmitLValue(L.IE).getPointer(CGF);
             Ptr = this->CGF.EmitLoadOfScalar(this->CGF.EmitLValue(L.IE),
                                              L.IE->getExprLoc());
-            // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the placeholder
-            // value MEMBER_OF=FFFF so that the entry is later updated with the
-            // correct value of MEMBER_OF.
+            // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the
+            // placeholder value MEMBER_OF=FFFF so that the entry is later
+            // updated with the correct value of MEMBER_OF.
             CurInfo.Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM |
                                     OMP_MAP_MEMBER_OF);
           }
           CurInfo.Exprs.push_back(L.VD);
           CurInfo.BasePointers.emplace_back(BasePtr, L.VD);
           CurInfo.Pointers.push_back(Ptr);
           CurInfo.Sizes.push_back(
               llvm::Constant::getNullValue(this->CGF.Int64Ty));
           CurInfo.Mappers.push_back(nullptr);
         }
       }
-
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
-      if (PartialStruct.Base.isValid())
+      if (PartialStruct.Base.isValid()) {
+        CurInfo.NonContigInfo.Dims.push_back(0);
         emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct, VD);
+      }
 
-      // We need to append the results of this capture to what we already have.
+      // We need to append the results of this capture to what we already
+      // have.
       CombinedInfo.append(CurInfo);
     }
     // Append data for use_device_ptr clauses.
     CombinedInfo.append(UseDevicePtrCombinedInfo);
   }
 
-  /// Generate all the base pointers, section pointers, sizes, map types, and
-  /// mappers for the extracted map clauses of user-defined mapper (all included
-  /// in \a CombinedInfo).
-  void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const {
-    assert(CurDir.is<const OMPDeclareMapperDecl *>() &&
-           "Expect a declare mapper directive");
-    const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
-    // We have to process the component lists that relate with the same
-    // declaration in a single chunk so that we can generate the map flags
-    // correctly. Therefore, we organize all lists in a map.
-    llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
-
-    // Fill the information map for map clauses.
-    for (const auto *C : CurMapperDir->clauselists()) {
-      const auto *MC = cast<OMPMapClause>(C);
-      const auto *EI = MC->getVarRefs().begin();
-      for (const auto L : MC->component_lists()) {
-        // The Expression is not correct if the mapping is implicit
-        const Expr *E = (MC->getMapLoc().isValid()) ? *EI : nullptr;
-        const ValueDecl *VD =
-            std::get<0>(L) ? cast<ValueDecl>(std::get<0>(L)->getCanonicalDecl())
-                           : nullptr;
-        // Get the corresponding user-defined mapper.
-        Info[VD].emplace_back(std::get<1>(L), MC->getMapType(),
-                              MC->getMapTypeModifiers(), llvm::None,
-                              /*ReturnDevicePointer=*/false, MC->isImplicit(),
-                              std::get<2>(L), E);
-        ++EI;
+public:
+  MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
+      : CurDir(&Dir), CGF(CGF) {
+    // Extract firstprivate clause information.
+    for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
+      for (const auto *D : C->varlists())
+        FirstPrivateDecls.try_emplace(
+            cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl()), C->isImplicit());
+    // Extract implicit firstprivates from uses_allocators clauses.
+    for (const auto *C : Dir.getClausesOfKind<OMPUsesAllocatorsClause>()) {
+      for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
+        OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
+        if (const auto *DRE = dyn_cast_or_null<DeclRefExpr>(D.AllocatorTraits))
+          FirstPrivateDecls.try_emplace(cast<VarDecl>(DRE->getDecl()),
+                                        /*Implicit=*/true);
+        else if (const auto *VD = dyn_cast<VarDecl>(
+                     cast<DeclRefExpr>(D.Allocator->IgnoreParenImpCasts())
+                         ->getDecl()))
+          FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
       }
     }
+    // Extract device pointer clause information.
+    for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
+      for (auto L : C->component_lists())
+        DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
+  }
 
-    for (const auto &M : Info) {
-      // We need to know when we generate information for the first component
-      // associated with a capture, because the mapping flags depend on it.
-      bool IsFirstComponentList = true;
-
-      // Underlying variable declaration used in the map clause.
-      const ValueDecl *VD = std::get<0>(M);
-
-      // Temporary generated information.
-      MapCombinedInfoTy CurInfo;
-      StructRangeInfoTy PartialStruct;
+  /// Constructor for the declare mapper directive.
+  MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
+      : CurDir(&Dir), CGF(CGF) {}
 
-      for (const MapInfo &L : M.second) {
-        assert(!L.Components.empty() &&
-               "Not expecting declaration with no component lists.");
-        generateInfoForComponentList(
-            L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo,
-            PartialStruct, IsFirstComponentList, L.IsImplicit, L.Mapper,
-            L.ForDeviceAddr, VD, L.VarRef);
-        IsFirstComponentList = false;
-      }
+  /// Generate code for the combined entry if we have a partially mapped struct
+  /// and take care of the mapping flags of the arguments corresponding to
+  /// individual struct members.
+  void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
+                         MapFlagsArrayTy &CurTypes,
+                         const StructRangeInfoTy &PartialStruct,
+                         const ValueDecl *VD = nullptr,
+                         bool NotTargetParams = true) const {
+    if (CurTypes.size() == 1 &&
+        ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) &&
+        !PartialStruct.IsArraySection)
+      return;
+    Address LBAddr = PartialStruct.LowestElem.second;
+    Address HBAddr = PartialStruct.HighestElem.second;
+    if (PartialStruct.HasCompleteRecord) {
+      LBAddr = PartialStruct.Base;
+      HBAddr = PartialStruct.Base;
+    CombinedInfo.Exprs.push_back(VD);
+    // Base is the base of the struct
+    CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer());
+    // Pointer is the address of the lowest element
+    llvm::Value *LB = LBAddr.getPointer();
+    CombinedInfo.Pointers.push_back(LB);
+    // There should not be a mapper for a combined entry.
+    CombinedInfo.Mappers.push_back(nullptr);
+    // Size is (addr of {highest+1} element) - (addr of lowest element)
+    llvm::Value *HB = HBAddr.getPointer();
+    llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1);
+    llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy);
+    llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy);
+    llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr);
+    llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty,
+                                                  /*isSigned=*/false);
+    CombinedInfo.Sizes.push_back(Size);
+    // Map type is always TARGET_PARAM, if generate info for captures.
+    CombinedInfo.Types.push_back(NotTargetParams ? OMP_MAP_NONE
+                                                 : OMP_MAP_TARGET_PARAM);
+    // If any element has the present modifier, then make sure the runtime
+    // doesn't attempt to allocate the struct.
+    if (CurTypes.end() !=
+        llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) {
+          return Type & OMP_MAP_PRESENT;
+        }))
+      CombinedInfo.Types.back() |= OMP_MAP_PRESENT;
+    // Remove TARGET_PARAM flag from the first element
+    (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM;
 
-      // If there is an entry in PartialStruct it means we have a struct with
-      // individual members mapped. Emit an extra combined entry.
-      if (PartialStruct.Base.isValid()) {
-        CurInfo.NonContigInfo.Dims.push_back(0);
-        emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct, VD);
-      }
+    // All other current entries will be MEMBER_OF the combined entry
+    // (except for PTR_AND_OBJ entries which do not have a placeholder value
+    // 0xFFFF in the MEMBER_OF field).
+    OpenMPOffloadMappingFlags MemberOfFlag =
+        getMemberOfFlag(CombinedInfo.BasePointers.size() - 1);
+    for (auto &M : CurTypes)
+      setCorrectMemberOfFlag(M, MemberOfFlag);
 
-      // We need to append the results of this capture to what we already have.
-      CombinedInfo.append(CurInfo);
-    }
+  /// Generate all the base pointers, section pointers, sizes, map types, and
+  /// mappers for the extracted mappable expressions (all included in \a
+  /// CombinedInfo). Also, for each item that relates with a device pointer, a
+  /// pair of the relevant declaration and index where it occurs is appended to
+  /// the device pointers info array.
+  void generateAllInfo(
+      MapCombinedInfoTy &CombinedInfo,
+      const llvm::DenseSet<CanonicalDeclPtr<const Decl>> &SkipVarSet =
+          llvm::DenseSet<CanonicalDeclPtr<const Decl>>()) const {
+    assert(CurDir.is<const OMPExecutableDirective *>() &&
+           "Expect a executable directive");
+    const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
+    generateAllInfoForClauses(CurExecDir->clauses(), CombinedInfo, SkipVarSet);
+
+  /// Generate all the base pointers, section pointers, sizes, map types, and
+  /// mappers for the extracted map clauses of user-defined mapper (all included
+  /// in \a CombinedInfo).
+  void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const {
+    assert(CurDir.is<const OMPDeclareMapperDecl *>() &&
+           "Expect a declare mapper directive");
+    const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
+    generateAllInfoForClauses(CurMapperDir->clauses(), CombinedInfo);
   }
 
   /// Emit capture info for lambdas for variables captured by reference.
   void generateInfoForLambdaCaptures(
       const ValueDecl *VD, llvm::Value *Arg, MapCombinedInfoTy &CombinedInfo,
       llvm::DenseMap<llvm::Value *, llvm::Value *> &LambdaPointers) const {
     const auto *RD = VD->getType()
                          .getCanonicalType()
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
         assert(!Components.empty() &&
                "Not expecting declaration with no component lists.");
         DeclComponentLists.emplace_back(Components, C->getMapType(),
                                         C->getMapTypeModifiers(),
                                         C->isImplicit(), Mapper, E);
         ++EI;
       }
     }
+    llvm::stable_sort(DeclComponentLists, [](const MapData &LHS,
+                                             const MapData &RHS) {
+      ArrayRef<OpenMPMapModifierKind> MapModifiers = std::get<2>(LHS);
+      OpenMPMapClauseKind MapType = std::get<1>(RHS);
+      bool HasPresent = !MapModifiers.empty() &&
+                        llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) {
+                          return K == clang::OMPC_MAP_MODIFIER_present;
+                        });
+      bool HasAllocs = MapType == OMPC_MAP_alloc;
+      MapModifiers = std::get<2>(RHS);
+      MapType = std::get<1>(LHS);
+      bool HasPresentR =
+          !MapModifiers.empty() &&
+          llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) {
+            return K == clang::OMPC_MAP_MODIFIER_present;
+          });
+      bool HasAllocsR = MapType == OMPC_MAP_alloc;
+      return (HasPresent && !HasPresentR) || (HasAllocs && !HasAllocsR);
+    });
 
     // Find overlapping elements (including the offset from the base element).
     llvm::SmallDenseMap<
         const MapData *,
         llvm::SmallVector<
             OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>,
         4>
         OverlappedData;
for (const MapData &L : DeclComponentLists) {
         for (; CI != CE && SI != SE; ++CI, ++SI) {
           if (CI->getAssociatedExpression()->getStmtClass() !=
               SI->getAssociatedExpression()->getStmtClass())
             break;
           // Are we dealing with different variables/fields?
           if (CI->getAssociatedDeclaration() != SI->getAssociatedDeclaration())
             break;
         }
-        // Found overlapping if, at least for one component, reached the head of
-        // the components list.
+        // Found overlapping if, at least for one component, reached the head
+        // of the components list.
         if (CI == CE || SI == SE) {
-          assert((CI != CE || SI != SE) &&
-                 "Unexpected full match of the mapping components.");
+          // Ignore it if it is the same component.
+          if (CI == CE && SI == SE)
+            continue;
+          const auto It = (SI == SE) ? CI : SI;
+          // If one component is a pointer and another one is a kind of
+          // dereference of this pointer (array subscript, section, dereference,
+          // etc.), it is not an overlapping.
+          if (!isa<MemberExpr>(It->getAssociatedExpression()) ||
+              std::prev(It)
+                  ->getAssociatedExpression()
+                  ->getType()
+                  .getNonReferenceType()
+                  ->isPointerType())
+            continue;
           const MapData &BaseData = CI == CE ? L : L1;
           OMPClauseMappableExprCommon::MappableExprComponentListRef SubData =
               SI == SE ? Components : Components1;
           auto &OverlappedElements = OverlappedData.FindAndConstruct(&BaseData);
           OverlappedElements.getSecond().push_back(SubData);
         }
       }
     }
     // Sort the overlapped elements for each item.
     llvm::SmallVector<const FieldDecl *, 4> Layout;
     if (!OverlappedData.empty()) {
       if (const auto *CRD =
               VD->getType().getCanonicalType()->getAsCXXRecordDecl())
         getPlainLayout(CRD, Layout, /*AsBase=*/false);
       else {
         const auto *RD = VD->getType().getCanonicalType()->getAsRecordDecl();
         Layout.append(RD->field_begin(), RD->field_end());
       }
     }
     for (auto &Pair : OverlappedData) {
-      llvm::sort(
+      llvm::stable_sort(
           Pair.getSecond(),
           [&Layout](
               OMPClauseMappableExprCommon::MappableExprComponentListRef First,
               OMPClauseMappableExprCommon::MappableExprComponentListRef
                   Second) {
             auto CI = First.rbegin();
             auto CE = First.rend();
             auto SI = Second.rbegin();
for (auto &Pair : OverlappedData) {
                   return FD == FD1 || FD == FD2;
                 });
             return *It == FD1;
           });
     }
 
     // Associated with a capture, because the mapping flags depend on it.
     // Go through all of the elements with the overlapped elements.
+    bool IsFirstComponentList = true;
     for (const auto &Pair : OverlappedData) {
       const MapData &L = *Pair.getFirst();
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
       const ValueDecl *Mapper;
       const Expr *VarRef;
       std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper, VarRef) =
           L;
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedComponents = Pair.getSecond();
-      bool IsFirstComponentList = true;
       generateInfoForComponentList(
           MapType, MapModifiers, llvm::None, Components, CombinedInfo,
           PartialStruct, IsFirstComponentList, IsImplicit, Mapper,
           /*ForDeviceAddr=*/false, VD, VarRef, OverlappedComponents);
+      IsFirstComponentList = false;
     }
     // Go through other elements without overlapped elements.
-    bool IsFirstComponentList = OverlappedData.empty();
     for (const MapData &L : DeclComponentLists) {
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
       const ValueDecl *Mapper;
       const Expr *VarRef;
       std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper, VarRef) =

 }
 
 /// Emit the user-defined mapper function. The code generation follows the
 /// pattern in the example below.
 /// \code
 /// void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 ///                                           void *base, void *begin,
 ///                                           int64_t size, int64_t type,
 ///                                           void *name = nullptr) {
-///   // Allocate space for an array section first.
-///   if (size > 1 && !maptype.IsDelete)
-///     __tgt_push_mapper_component(rt_mapper_handle, base, begin,
-///                                 size*sizeof(Ty), clearToFrom(type));
 ///   // Map members.
 ///   for (unsigned i = 0; i < size; i++) {
 ///     // For each component specified by this mapper:
 ///     for (auto c : all_components) {
 ///       if (c.hasMapper())
 ///         (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size,
 ///                       c.arg_type, c.arg_name);
 ///       else
 ///         __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
 ///                                     c.arg_begin, c.arg_size, c.arg_type,
 ///                                     c.arg_name);
-///     }
-///   }
-///   // Delete the array section.
-///   if (size > 1 && maptype.IsDelete)
+///   // Delete/allocate space for an array section first.
+///   if (size > 1)
 ///     __tgt_push_mapper_component(rt_mapper_handle, base, begin,
 ///                                 size*sizeof(Ty), clearToFrom(type));
-/// }
+///     }
+///   }
+/// }
 /// \endcode
 void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
                                             CodeGenFunction *CGF) {
   if (UDMMap.count(D) > 0)
     return;
   ASTContext &C = CGM.getContext();
   QualType Ty = D->getType();
   QualType PtrTy = C.getPointerType(Ty).withRestrict();
llvm::Value *Handle = MapperCGF.EmitLoadOfScalar(
       /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
   llvm::Value *BaseIn = MapperCGF.EmitLoadOfScalar(
       MapperCGF.GetAddrOfLocalVar(&BaseArg),
       /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
   llvm::Value *BeginIn = MapperCGF.EmitLoadOfScalar(
       MapperCGF.GetAddrOfLocalVar(&BeginArg),
       /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
 
-  // Emit array initiation if this is an array section and \p MapType indicates
-  // that memory allocation is required.
-  llvm::BasicBlock *HeadBB = MapperCGF.createBasicBlock("omp.arraymap.head");
-  emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType,
-                             ElementSize, HeadBB, /*IsInit=*/true);
-
-  // Emit a for loop to iterate through SizeArg of elements and map all of them.
-
-  // Emit the loop header block.
-  MapperCGF.EmitBlock(HeadBB);
   llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.arraymap.body");
   llvm::BasicBlock *DoneBB = MapperCGF.createBasicBlock("omp.done");
   // Evaluate whether the initial condition is satisfied.
   llvm::Value *IsEmpty =
       MapperCGF.Builder.CreateICmpEQ(PtrBegin, PtrEnd, "omp.arraymap.isempty");
   MapperCGF.Builder.CreateCondBr(IsEmpty, DoneBB, BodyBB);
   llvm::BasicBlock *EntryBB = MapperCGF.Builder.GetInsertBlock();
 
llvm::BasicBlock *MemberCombineBB =
         MapperCGF.createBasicBlock("omp.member.combine");
     llvm::BasicBlock *TypeBB = MapperCGF.createBasicBlock("omp.type");
     llvm::Value *IsMember = MapperCGF.Builder.CreateIsNull(Member);
     MapperCGF.Builder.CreateCondBr(IsMember, TypeBB, MemberCombineBB);
     // Add the number of pre-existing components to the MEMBER_OF field if it
     // is valid.
     MapperCGF.EmitBlock(MemberCombineBB);
     llvm::Value *CombinedMember =
         MapperCGF.Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
     // Do nothing if it is not a member of previous components.
     MapperCGF.EmitBlock(TypeBB);
     llvm::PHINode *MemberMapType =
         MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.membermaptype");
     MemberMapType->addIncoming(OriMapType, MemberBB);
     MemberMapType->addIncoming(CombinedMember, MemberCombineBB);
 
     // Combine the map type inherited from user-defined mapper with that
llvm::Value *IsDone =
       MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone");
   llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit");
   MapperCGF.Builder.CreateCondBr(IsDone, ExitBB, BodyBB);
 
   MapperCGF.EmitBlock(ExitBB);
   // Emit array deletion if this is an array section and \p MapType indicates
   // that deletion is required.
   emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType,
-                             ElementSize, DoneBB, /*IsInit=*/false);
+                             ElementSize, DoneBB);
 
   // Emit the function exit block.
   MapperCGF.EmitBlock(DoneBB, /*IsFinished=*/true);
   MapperCGF.FinishFunction();
   UDMMap.try_emplace(D, Fn);
   if (CGF) {
     auto &Decls = FunctionUDMMap.FindAndConstruct(CGF->CurFn);
     Decls.second.push_back(D);
   }
 }
 
 /// Emit the array initialization or deletion portion for user-defined mapper
 /// code generation. First, it evaluates whether an array section is mapped and
 /// whether the \a MapType instructs to delete this section. If \a IsInit is
 /// true, and \a MapType indicates to not delete this array, array
 /// initialization code is generated. If \a IsInit is false, and \a MapType
 /// indicates to not this array, array deletion code is generated.
 void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
     CodeGenFunction &MapperCGF, llvm::Value *Handle, llvm::Value *Base,
     llvm::Value *Begin, llvm::Value *Size, llvm::Value *MapType,
-    CharUnits ElementSize, llvm::BasicBlock *ExitBB, bool IsInit) {
-  StringRef Prefix = IsInit ? ".init" : ".del";
-
+    CharUnits ElementSize, llvm::BasicBlock *ExitBB) {
   // Evaluate if this is an array section.
-  llvm::BasicBlock *IsDeleteBB =
-      MapperCGF.createBasicBlock(getName({"omp.array", Prefix, ".evaldelete"}));
-  llvm::BasicBlock *BodyBB =
-      MapperCGF.createBasicBlock(getName({"omp.array", Prefix}));
+  llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.array.body");
   llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE(
       Size, MapperCGF.Builder.getInt64(1), "omp.arrayinit.isarray");
-  MapperCGF.Builder.CreateCondBr(IsArray, IsDeleteBB, ExitBB);
-
-  // Evaluate if we are going to delete this section.
-  MapperCGF.EmitBlock(IsDeleteBB);
-  llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd(
-      MapType,
-      MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE));
-  llvm::Value *DeleteCond;
-  if (IsInit) {
-    DeleteCond = MapperCGF.Builder.CreateIsNull(
-        DeleteBit, getName({"omp.array", Prefix, ".delete"}));
-  } else {
-    DeleteCond = MapperCGF.Builder.CreateIsNotNull(
-        DeleteBit, getName({"omp.array", Prefix, ".delete"}));
-  }
-  MapperCGF.Builder.CreateCondBr(DeleteCond, BodyBB, ExitBB);
+  MapperCGF.Builder.CreateCondBr(IsArray, BodyBB, ExitBB);
 
   MapperCGF.EmitBlock(BodyBB);
   // Get the array size by multiplying element size and element number (i.e., \p
   // Size).
   llvm::Value *ArraySize = MapperCGF.Builder.CreateNUWMul(
       Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity()));
   // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves
   // memory allocation/deletion purpose only.
auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
     MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
 
     // Get mappable expression information.
     MappableExprsHandler MEHandler(D, CGF);
     llvm::DenseMap<llvm::Value *, llvm::Value *> LambdaPointers;
     llvm::DenseSet<CanonicalDeclPtr<const Decl>> MappedVarSet;
 
     auto RI = CS.getCapturedRecordDecl()->field_begin();
-    auto CV = CapturedVars.begin();
+    auto *CV = CapturedVars.begin();
     for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
                                               CE = CS.capture_end();
          CI != CE; ++CI, ++RI, ++CV) {
       MappableExprsHandler::MapCombinedInfoTy CurInfo;
       MappableExprsHandler::StructRangeInfoTy PartialStruct;
 
       // VLA sizes are passed to the outlined region by copy and do not have map
       // information associated.

clang/lib/Sema/SemaOpenMP.cpp

if (auto *VD = dyn_cast<VarDecl>(E->getDecl())) {
           }
         }
       }
 
       if (isOpenMPTargetExecutionDirective(DKind) &&
           !Stack->isLoopControlVariable(VD).first) {
         if (!Stack->checkMappableExprComponentListsForDecl(
                 VD, /*CurrentRegionOnly=*/true,
-                [](OMPClauseMappableExprCommon::MappableExprComponentListRef
+                [this](OMPClauseMappableExprCommon::MappableExprComponentListRef
                            StackComponents,
                        OpenMPClauseKind) {
+                  if (SemaRef.LangOpts.OpenMP >= 50)
+                    return !StackComponents.empty();
                   // Variable is used if it has been marked as an array, array
                   // section, array shaping or the variable iself.
                   return StackComponents.size() == 1 ||
                          std::all_of(
                              std::next(StackComponents.rbegin()),
                              StackComponents.rend(),
                              [](const OMPClauseMappableExprCommon::
                                     MappableComponent &MC) {
static bool checkMapConflicts(
   const Expr *EnclosingExpr = nullptr;
 
   bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
       VD, CurrentRegionOnly,
       [&IsEnclosedByDataEnvironmentExpr, &SemaRef, VD, CurrentRegionOnly, ELoc,
        ERange, CKind, &EnclosingExpr,
        CurComponents](OMPClauseMappableExprCommon::MappableExprComponentListRef
                           StackComponents,
-                      OpenMPClauseKind) {
+                      OpenMPClauseKind Kind) {
+        if (CKind == Kind && SemaRef.LangOpts.OpenMP >= 50)
+          return false;
         assert(!StackComponents.empty() &&
                "Map clause expression with no components!");
         assert(StackComponents.back().getAssociatedDeclaration() == VD &&
                "Map clause expression with unexpected base!");
         (void)VD;
 
         // The whole expression in the stack.
         const Expr *RE = StackComponents.front().getAssociatedExpression();
for (Expr *RE : MVLI.VarList) {
     // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4]
     // A list item in a to or from clause must have a mappable type.
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
     //  A list item must have a mappable type.
     if (!checkTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
                            DSAS, Type))
       continue;
 
-    Type = I->getAssociatedDeclaration()->getType().getNonReferenceType();
-
     if (CKind == OMPC_map) {
       // target enter data
       // OpenMP [2.10.2, Restrictions, p. 99]
       // A map-type must be specified in all map clauses and must be either
       // to or alloc.
       OpenMPDirectiveKind DKind = DSAS->getCurrentDirective();
       if (DKind == OMPD_target_enter_data &&
           !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {

clang/test/OpenMP/declare_mapper_codegen.cpp

 // CK0-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
 // CK0-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
 // CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
 // CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
 // CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
 // CK0-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
 // CK0-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
-// CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
-// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
-
-// CK0: [[INITEVALDEL]]
-// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
-// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
-// CK0: [[INIT]]
-// CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
-// CK0: br label %[[LHEAD:[^,]+]]
-
-// CK0: [[LHEAD]]
 // CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
 // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK0: [[LBODY]]
-// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
 // CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
 // CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
 // CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
 // CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
 // CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
 // CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
 // CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*

 // CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
 // CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
 // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // CK0: [[LEXIT]]
 // CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
 // CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK0: [[EVALDEL]]
-// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK0: [[DEL]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
 // CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
 // CK0: br label %[[DONE]]
 // CK0: [[DONE]]
 // CK0: ret void
 

 // CK1-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
 // CK1-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 4
 // CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
 // CK1-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
 // CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
 // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
 // CK1-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
 // CK1-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
-// CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
-// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
-
-// CK1: [[INITEVALDEL]]
-// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
-// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
-// CK1: [[INIT]]
-// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
-// CK1: br label %[[LHEAD:[^,]+]]
-
-// CK1: [[LHEAD]]
 // CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
 // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK1: [[LBODY]]
-// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
 // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
 // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
 // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 // CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
 // CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
 // CK1-DAG: br label %[[MEMBER:[^,]+]]
 // CK1-DAG: [[MEMBER]]
 // CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
 // CK1-DAG: [[MEMBERCOM]]
-// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]]
+// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
 // CK1-DAG: br label %[[LTYPE]]
 // CK1-DAG: [[LTYPE]]
-// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 3, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 // CK1-DAG: [[ALLOC]]
 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 // CK1-DAG: br label %[[TYEND:[^,]+]]
 // CK1-DAG: [[ALLOCELSE]]
 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1

 // CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
 // CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // CK1: [[LEXIT]]
 // CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
 // CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK1: [[EVALDEL]]
-// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK1: [[DEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
 // CK1: br label %[[DONE]]
 // CK1: [[DONE]]
 // CK1: ret void
 
 #endif // CK1

 // CK2-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
 // CK2-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
 // CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
 // CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
 // CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
 // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
 // CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
 // CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
-// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
-// CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
-
-// CK2: [[INITEVALDEL]]
-// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
-// CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
-// CK2: [[INIT]]
-// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
-// CK2: br label %[[LHEAD:[^,]+]]
-
-// CK2: [[LHEAD]]
 // CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
 // CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK2: [[LBODY]]
-// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
 // CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
 // CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
 // CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 // CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
 // CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
 // CK2-DAG: br label %[[MEMBER:[^,]+]]
 // CK2-DAG: [[MEMBER]]
 // CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
 // CK2-DAG: [[MEMBERCOM]]
-// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]]
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
 // CK2-DAG: br label %[[LTYPE]]
 // CK2-DAG: [[LTYPE]]
-// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 3, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
 // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 // CK2-DAG: [[ALLOC]]
 // CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 // CK2-DAG: br label %[[TYEND:[^,]+]]
 // CK2-DAG: [[ALLOCELSE]]
 // CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1

 // CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
 // CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
 // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // CK2: [[LEXIT]]
 // CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
 // CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK2: [[EVALDEL]]
-// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK2: [[DEL]]
 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 // CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
 // CK2: br label %[[DONE]]
 // CK2: [[DONE]]
 // CK2: ret void
 
 #endif // CK2

 // CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
 // CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
 // CK4-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
 // CK4-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
 // CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
 // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
 // CK4-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
 // CK4-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
-// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
-// CK4: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
-
-// CK4: [[INITEVALDEL]]
-// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK4: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
-// CK4: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
-// CK4: [[INIT]]
-// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
-// CK4: br label %[[LHEAD:[^,]+]]
-
-// CK4: [[LHEAD]]
 // CK4: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
 // CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 // CK4: [[LBODY]]
-// CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 // CK4: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
 // CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
 // CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
 // CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
 // CK4-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
 // CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
 // CK4-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
 // CK4-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*

 // CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
 // CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
 // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // CK4: [[LEXIT]]
 // CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
 // CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 // CK4: [[EVALDEL]]
-// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK4: [[DEL]]
 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
 // CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
 // CK4: br label %[[DONE]]
 // CK4: [[DONE]]
 // CK4: ret void
 

clang/test/OpenMP/target_data_codegen.cpp

 };
 struct S2 {
   S1 s;
   struct S2 *ps;
 };
 
 void test_close_modifier(int arg) {
   S2 *ps;
-  // CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
-  #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
+#pragma omp target data map(close, tofrom \
+                            : arg, ps->ps->ps->ps->s)
   {
     ++(arg);
   }
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64
 // RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-64
 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
 // RUN: %clang_cc1 -DCK6 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
 
 // RUN: %clang_cc1 -DCK6 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK6 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 #ifdef CK6
 void test_close_modifier(int arg) {
-  // CK6: private unnamed_addr constant [1 x i64] [i64 1027]
-  #pragma omp target data map(close,tofrom: arg)
+#pragma omp target data map(close, tofrom \
+                            : arg)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
 // RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7 --check-prefix CK7-64
 
void test_present_modifier(int arg) {
   S2 *ps1;
   S2 *ps2;
 
   // Make sure the struct picks up present even if another element of the struct
   // doesn't have present.
 
   // CK8: private unnamed_addr constant [11 x i64]
 
-  // ps1
-  //
-  // PRESENT=0x1000 = 0x1000
-  // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
-  // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
-  // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
-  // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
-  //
-  // CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000000003]],
-  // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],
+// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],
 
-  // arg
-  //
-  // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
-  //
-  // CK8-SAME: {{^}} i64 [[#0x1003]],
 
-  // ps2
-  //
-  // PRESENT=0x1000 = 0x1000
-  // MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
-  // MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
-  // PTR_AND_OBJ=0x10 = 0x10
-  // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
-  //
-  // CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
-  // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
-  #pragma omp target data map(tofrom: ps1->s) \
-                          map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
-                          map(tofrom: ps2->ps->ps->ps->s)
+#pragma omp target data map(tofrom         \
+                            : ps1->s)      \
+    map(present, tofrom                    \
+        : arg, ps1->ps->ps->ps->s, ps2->s) \
+        map(tofrom                         \
+            : ps2->ps->ps->ps->s)
   {
     ++(arg);
   }
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64
 // RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-64
 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
 // RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
 
 // RUN: %clang_cc1 -DCK9 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK9 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 #ifdef CK9
 void test_present_modifier(int arg) {
-  // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
-  // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]
-  #pragma omp target data map(present,tofrom: arg)
+#pragma omp target data map(present, tofrom \
+                            : arg)
   {++arg;}
 }
 #endif
 #endif

clang/test/OpenMP/target_enter_data_codegen.cpp

 };
 
 ST<int> gb;
 double gc[100];
 
 // PRESENT=0x1000 | TO=0x1 = 0x1001
 // CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
 
-// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1425
+// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405
 // CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
 
 // CK1A-LABEL: _Z3fooi
 void foo(int arg) {
   int la;
   float lb[arg];
 
   // Region 00

clang/test/OpenMP/target_is_device_ptr_messages.cpp

-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -ferror-limit 200 %s -Wuninitialized
 
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
 struct ST {
   int *a;
 };
 typedef int arr[10];
 typedef ST STarr[10];
 struct SA {
   const int d = 5;
   const int da[5] = { 0 };

clang/test/OpenMP/target_map_codegen_29.cpp

 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
 // CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i64]* [[SIZES]], i32 0, i32 0
 // CK30-DAG: store i64 [[S_ALLOC_SIZE:%.+]], i64* [[SIZE]],
 // CK30-DAG: [[S_ALLOC_SIZE]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_BEGIN_BC:%.+]]
 // CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
 // CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
 // CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
-// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i32 1
-// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
-// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
+// CK30-DAG: [[S_END]] = bitcast [[STRUCT]]* [[REAL_S_END:%.+]] to i8*
+// CK30-DAG: [[REAL_S_END]] = getelementptr [[STRUCT]], [[STRUCT]]* [[S]], i32 1
 
 // CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
 // CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 1
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
 // CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 1

 // CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 3
 // CK30-DAG: store i64 [[SIZE3:%.+]], i64* [[SIZE]],
 // CK30-DAG: [[PTR2]] = getelementptr i32*, i32** [[S_PTR1_REF]], i{{64|32}} 1
 // CK30-DAG: [[SIZE3]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_PTR1_BC:%.+]]
 // CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
 // CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
 // CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8*
-// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST]], i{{64|32}} 1
+// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{64|32}} 1
+// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
+// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
 
 // CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
 // CK30-DAG: store i32** [[S_PTR1:%.+]], i32*** [[BC]],
 // CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 4
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32**
 // CK30-DAG: store i32* [[S_PTR1_BEGIN:%.+]], i32** [[BC]],
 // CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 4

clang/test/OpenMP/target_map_codegen_31.cpp

 };
 
 // CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 //
 // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
 // CK31A-USE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]],
 // CK31A-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1000]],
 //
-// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
 // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
-// CK31A-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1023]],
-// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1003]],
+// CK31A-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1023]],
+// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1003]],
 //
 // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
 // MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x5000000001003
 // MEMBER_OF_5=0x5000000000000 | FROM=0x2 | TO=0x1 = 0x5000000000003
 // CK31A-USE-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x5000000001003]], i64 [[#0x5000000000003]]]
 // CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x5000000001003]], i64 [[#0x5000000000003]]]
 
 // CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
void explicit_maps_single (int ii){
   // CK31A: [[A:%.+]] = alloca i32
   int a = ii;
 
   // CK31A: [[ST1:%.+]] = alloca [[ST]]
   // CK31A: [[ST2:%.+]] = alloca [[ST]]
   struct ST st1;
   struct ST st2;
 
-  // Make sure the struct picks up present even if another element of the struct
-  // doesn't have present.
-  // Region 00
-  // CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
-  // CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
-  // CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
-  // CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
-  // CK31A-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
-  // CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
-  // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-  // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
 
-  // st1
-  // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-  // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-  // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-  // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
-  // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
-  // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
 
-  // st1.i
-  // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-  // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-  // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
-  // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
-  // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP1]]
-  // CK31A-DAG: store i64 4, i64* [[S1]]
+// st1.j
+// CK31A-DAG: store i32* [[ST1_J]], i32** [[CP1]]
 
-  // st1.j
-  // CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
-  // CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
-  // CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
-  // CK31A-DAG: store i32* [[ST1_J]], i32** [[CP2]]
-  // CK31A-DAG: store i64 4, i64* [[S2]]
+// st1.i
+// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP2]]
 
-  // a
-  // CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
-  // CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
-  // CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
-  // CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
-  // CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
-  // CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
-  // CK31A-DAG: store i32* [[A]], i32** [[CP3]]
-  // CK31A-DAG: store i64 4, i64* [[S3]]
 
-  // st2
-  // CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
-  // CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
-  // CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
-  // CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
-  // CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
-  // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
-  // CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
 
-  // st2.i
-  // CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
-  // CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
-  // CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
-  // CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
-  // CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
-  // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
-  // CK31A-DAG: store i64 4, i64* [[S5]]
 
-  // st2.j
-  // CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
-  // CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
-  // CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
-  // CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
-  // CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
-  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
-  // CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
-  // CK31A-DAG: store i64 4, i64* [[S6]]
 
-  // CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
-  // CK31A-NOUSE: call void [[CALL00:@.+]]()
-  #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
+#pragma omp target map(tofrom                                     \
+                       : st1.i) map(present, tofrom               \
+                                    : a, st1.j, st2.i) map(tofrom \
+                                                           : st2.j)
   {
 #ifdef USE
     st1.i++;
     a++;
     st1.j++;
     st2.i++;
     st2.j++;
 #endif

clang/test/OpenMP/target_map_codegen_32.cpp

 // RUN: %clang_cc1 -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
 #ifdef CK31B
 
 // CK31B: [[ST:%.+]] = type { i32, i32 }
 
 // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
-// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
 
 // CK31B-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK31B-USE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1020]],
 // CK31B-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1000]],
-// CK31B-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]]
-// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]]
+// CK31B-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]]
+// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]]
 
 struct ST {
   int i;
   int j;
   // CK31B-LABEL: define {{.*}}test_present_members{{.*}}(
   void test_present_members() {
-    // Make sure the struct picks up present even if another element of the
-    // struct doesn't have present.
-    // Region 00
-    // CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0
-    // CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
-    // CK31B-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}})
-    // CK31B-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
-    // CK31B-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-    // CK31B-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+// CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 1
+// CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 0
 
-    // this
-    // CK31B-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-    // CK31B-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-    // CK31B-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-    // CK31B-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
-    // CK31B-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]]
-    // CK31B-DAG: store i32* [[I]], i32** [[CP0]]
-    // CK31B-DAG: store i64 %{{.+}}, i64* [[S0]]
 
-    // i
-    // CK31B-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-    // CK31B-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-    // CK31B-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-    // CK31B-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
-    // CK31B-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]]
-    // CK31B-DAG: store i32* [[I]], i32** [[CP1]]
-    // CK31B-DAG: store i64 4, i64* [[S1]]
+// j
+// CK31B-DAG: store i32* [[J]], i32** [[CP1]]
 
-    // j
-    // CK31B-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-    // CK31B-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-    // CK31B-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
-    // CK31B-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
-    // CK31B-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
-    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]]
-    // CK31B-DAG: store i32* [[J]], i32** [[CP2]]
-    // CK31B-DAG: store i64 4, i64* [[S2]]
+// i
+// CK31B-DAG: store i32* [[I]], i32** [[CP2]]
 
-    // CK31B-USE: call void [[CALL00:@.+]]([[ST]]* [[THIS]])
-    // CK31B-NOUSE: call void [[CALL00:@.+]]()
-    #pragma omp target map(tofrom: i) map(present, tofrom: j)
+#pragma omp target map(tofrom                   \
+                       : i) map(present, tofrom \
+                                : j)
     {
 #ifdef USE
       i++;
       j++;
 #endif
     }
   }
 };

clang/test/OpenMP/target_map_messages.cpp

void func(int arg) {
     #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
     {}
 
     #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
     {}
 
     #pragma omp target map(to:b,e)
     {}
-    #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+    #pragma omp target map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
     {}
     #pragma omp target map(to:b[:2],e)
     {}
     #pragma omp target map(to:b,e[:])
     {}
     #pragma omp target map(b[-1:]) // expected-error {{array section must be a subset of the original array}}
     {}
     #pragma omp target map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}}

 #pragma omp target map(r.ArrS[0].Arr [1:23])
   {}
 #pragma omp target map(r.ArrS[0].Arr [1:arg])
   {}
 #pragma omp target map(r.ArrS[0].Arr [arg:23])
   {}
 #pragma omp target map(r.ArrS[0].Error)          // expected-error {{no member named 'Error' in 'SB'}}
   {}
-#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}}
   {}
 #pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
   {}
-#pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
   {}
-#pragma omp target map(r.PtrS, r.PtrS->B)    // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.PtrS, r.PtrS->B)    // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   {}
 #pragma omp target map(r.PtrS->A, r.PtrS->B)
   {}
-#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
   {}
-#pragma omp target map(r.RPtrS, r.RPtrS->B)    // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS, r.RPtrS->B)    // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   {}
 #pragma omp target map(r.RPtrS->A, r.RPtrS->B)
   {}
 #pragma omp target map(r.S.Arr[:12])
   {}
 #pragma omp target map(r.S.foo() [:12]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
   {}
 #pragma omp target map(r.C, r.D)
   {}
-#pragma omp target map(r.C, r.C)     // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C, r.C)     // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
-#pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
 #pragma omp target map(r.C, r.S)     // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
   {}
-#pragma omp target map(r, r.S)       // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r, r.S)       // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
 #pragma omp target map(r.C, t.C)
   {}
 #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
   {}
 #pragma omp target map(r.Arr)
   {}
 #pragma omp target map(r.Arr [3:5])
#pragma omp target map(u.B)        // expected-error {{mapping of union members is not allowed}}
   {}
 #pragma omp target
   {
     u.B = 0;
     r.S.foo();
   }
 
 #pragma omp target data map(to \
-                            : r.C) //expected-note {{used here}}
+                            : r.C) // lt50-note {{used here}}
   {
-#pragma omp target map(r.D)        // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+#pragma omp target map(r.D)        // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
     {}
   }
 
 #pragma omp target data map(to \
-                            : t.Ptr) //expected-note {{used here}}
+                            : t.Ptr) // lt50-note {{used here}}
   {
-#pragma omp target map(t.Ptr[:23])   // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target map(t.Ptr[:23])   // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
     {}
   }
 
 #pragma omp target data map(to \
                             : t.C, t.D)
   {
 #pragma omp target data map(to \
                             : t.C)

 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning 2 {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ba) // warn-warning 2 {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning 2 {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g) // warn-warning 2 {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
+#pragma omp target map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   foo();
 #pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
-#pragma omp target map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}

 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba) // warn-warning {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g) // warn-warning {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+#pragma omp target map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   foo();
 #pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
-#pragma omp target map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}

 #pragma omp target data map(tofrom, present: x)
   foo();
 #pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}  expected-note {{defined as private}}
   {}
 #pragma omp target firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
   {}
 #pragma omp target map(m) // warn-warning {{Type 'S6<int>' is not trivially copyable and not guaranteed to be mapped correctly}}
   {}
-// expected-note@+1 {{used here}}
+#pragma omp target
+  { s.a++; }
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.s)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.s.a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.b[:5])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
 #pragma omp target map(s.p[:5])
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.sa[3].a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.sp[3]->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.s.b[:2])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.s.p->b[:2])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
 #pragma omp target map(s.p->p->p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
 #pragma omp target map(s.s.s.b[:2])
   { s.s.s.b[0]++; }
 #pragma omp target map(s8[0:1], s9) // warn-warning {{Type 'class S8' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'class S9' is not trivially copyable and not guaranteed to be mapped correctly}}
   {}
 
   int **BB, *offset, *a;
 

clang/test/OpenMP/target_parallel_for_map_messages.cpp

 #pragma omp target parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)       // lt50-note 2 {{used here}}
+#pragma omp target parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

 #pragma omp target parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)       // lt50-note {{used here}}
+#pragma omp target parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp

 #pragma omp target parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)            // lt50-note 2 {{used here}}
+#pragma omp target parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

 #pragma omp target parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)            // lt50-note {{used here}}
+#pragma omp target parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

clang/test/OpenMP/target_parallel_map_messages.cpp

 #pragma omp target parallel map(S2::S2s)
   foo();
 #pragma omp target parallel map(S2::S2sc)
   foo();
 #pragma omp target parallel map(e, g)
   foo();
 #pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   foo();
-#pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   foo();
-#pragma omp target parallel map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   foo();
 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(da[:4])
   foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)   // lt50-note 2 {{used here}}
+#pragma omp target parallel map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   foo();
 #pragma omp target parallel map(j)
   foo();
-#pragma omp target parallel map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target parallel map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
     foo();
 #pragma omp target parallel map(j)
   foo();
 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();

 #pragma omp target parallel map(S2::S2s)
   foo();
 #pragma omp target parallel map(S2::S2sc)
   foo();
 #pragma omp target parallel map(e, g)
   foo();
 #pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   foo();
-#pragma omp target parallel map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   foo();
-#pragma omp target parallel map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   foo();
 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(da[:4])
   foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)   // lt50-note {{used here}}
+#pragma omp target parallel map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   foo();
 #pragma omp target parallel map(j)
   foo();
-#pragma omp target parallel map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target parallel map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target parallel map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
     foo();
 #pragma omp target parallel map(j)
   foo();
 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();

clang/test/OpenMP/target_simd_map_messages.cpp

 #pragma omp target simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target simd map(k[:4])   // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(l) // OK
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(da[:4])
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target simd map(k[:4])   // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(j)
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(l) //
   for (i = 0; i < argc; ++i) foo();
   }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
 // expected-error@+1 {{missing map type}}

clang/test/OpenMP/target_teams_distribute_map_messages.cpp

 #pragma omp target teams distribute map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)           // lt50-note 2 {{used here}}
+#pragma omp target teams distribute map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target teams distribute map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

 #pragma omp target teams distribute map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)           // lt50-note {{used here}}
+#pragma omp target teams distribute map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target teams distribute map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp

 #pragma omp target teams distribute parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)                        // lt50-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target teams distribute parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(da[:4])
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target data map(k, j, l)                        // lt50-note {{used here}}
+#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(j)
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
-  for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+  for (i = 0; i < argc; ++i)
+    foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target teams distribute parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute parallel for map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
   }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
 // expected-error@+1 {{missing map type}}

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp

 #pragma omp target teams distribute parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)                             // lt50-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target teams distribute parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

 #pragma omp target teams distribute parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)                             // lt50-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target teams distribute parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute parallel for simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp

 #pragma omp target teams distribute simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)                // lt50-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
   {
-#pragma omp target teams distribute simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

 #pragma omp target teams distribute simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(e, g)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(da[:4])
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l)                // lt50-note {{used here}}
+#pragma omp target teams distribute simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(j)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
+#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
   {
-#pragma omp target teams distribute simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
-  for (i = 0; i < argc; ++i) foo();
+#pragma omp target teams distribute simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+    for (i = 0; i < argc; ++i)
+      foo();
 #pragma omp target teams distribute simd map(j)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();

clang/test/OpenMP/target_teams_map_messages.cpp

void func(int arg) {
     #pragma omp target teams map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
     {}
 
     #pragma omp target teams map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
     {}
 
     #pragma omp target teams map(to:b,e)
     {}
-    #pragma omp target teams map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+    #pragma omp target teams map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
     {}
     #pragma omp target teams map(to:b[:2],e)
     {}
     #pragma omp target teams map(to:b,e[:])
     {}
     #pragma omp target teams map(b[-1:]) // expected-error {{array section must be a subset of the original array}}
     {}
     #pragma omp target teams map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}}
void SAclient(int arg) {
   #pragma omp target teams map(r.ArrS[0].Arr[1:23])
   {}
   #pragma omp target teams map(r.ArrS[0].Arr[1:arg])
   {}
   #pragma omp target teams map(r.ArrS[0].Arr[arg:23])
   {}
   #pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
   {}
-  #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+  #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}}
   {}
   #pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A)
   {}
-  #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
   {}
-  #pragma omp target teams map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+  #pragma omp target teams map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   {}
   #pragma omp target teams map(r.PtrS->A, r.PtrS->B)
   {}
-  #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
   {}
-  #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+  #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   {}
   #pragma omp target teams map(r.RPtrS->A, r.RPtrS->B)
   {}
   #pragma omp target teams map(r.S.Arr[:12])
   {}
   // ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
   // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
   #pragma omp target teams map(r.S.foo()[:12])
   {}
   #pragma omp target teams map(r.C, r.D)
   {}
-  #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  #pragma omp target teams map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
-  #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  #pragma omp target teams map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
   #pragma omp target teams map(r.C, r.S)  // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
   {}
-  #pragma omp target teams map(r, r.S)  // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  #pragma omp target teams map(r, r.S)  // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   {}
   #pragma omp target teams map(r.C, t.C)
   {}
   #pragma omp target teams map(r.A)   // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
   {}
   #pragma omp target teams map(r.Arr)
   {}
   #pragma omp target teams map(r.Arr[3:5])
void SAclient(int arg) {
   #pragma omp target teams map(r.S.Ptr[:])  // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
   {}
   // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
   #pragma omp target teams map((p+1)->A)
   {}
   #pragma omp target teams map(u.B)  // expected-error {{mapping of union members is not allowed}}
   {}
 
-  #pragma omp target data map(to: r.C) //expected-note {{used here}}
+  #pragma omp target data map(to: r.C) // lt50-note {{used here}}
   {
-    #pragma omp target teams map(r.D)  // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+    #pragma omp target teams map(r.D)  // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
     {}
   }
 
-  #pragma omp target data map(to: t.Ptr) //expected-note {{used here}}
+  #pragma omp target data map(to: t.Ptr) // lt50-note {{used here}}
   {
-    #pragma omp target teams map(t.Ptr[:23])  // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+    #pragma omp target teams map(t.Ptr[:23])  // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
     {}
   }
 
   #pragma omp target data map(to: t.C, t.D)
   {
   #pragma omp target data map(to: t.C)
   {
     #pragma omp target teams map(t.D)

 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #pragma omp target data map(ba)
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g)
 #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
-#pragma omp target teams map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
   foo();
 
 #pragma omp target data map(da)
 #pragma omp target teams map(da[:4])
   foo();
 
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
-#pragma omp target teams map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
   foo();
 
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}

 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba)
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g)
 #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-#pragma omp target teams map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
   foo();
 
 #pragma omp target data map(da)
 #pragma omp target teams map(da[:4])
   foo();
 
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
-#pragma omp target teams map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
   foo();
 
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
 // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}

clang/test/OpenMP/target_update_codegen.cpp

 // RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK19 -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK19
 
-// PRESENT=0x1000 | TO=0x1 = 0x1021
+// PRESENT=0x1000 | TO=0x1 = 0x1001
 // CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
 
 // PRESENT=0x1000 | FROM=0x2 = 0x1002
 // CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
 
 // CK19-LABEL: _Z13check_presenti
 void check_present(int arg) {
   int la;

openmp/libomptarget/src/omptarget.cpp

int targetDataMapper(ident_t *loc, DeviceTy &Device, void *arg_base, void *arg,
   std::vector<void *> MapperArgs(MapperComponents.Components.size());
   std::vector<int64_t> MapperArgSizes(MapperComponents.Components.size());
   std::vector<int64_t> MapperArgTypes(MapperComponents.Components.size());
   std::vector<void *> MapperArgNames(MapperComponents.Components.size());
 
   for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {
     auto &C =
         MapperComponents
-            .Components[target_data_function == targetDataEnd ? I : E - I - 1];
+            .Components[target_data_function == targetDataEnd ? E - I - 1 : I];
     MapperArgsBase[I] = C.Base;
     MapperArgs[I] = C.Begin;
     MapperArgSizes[I] = C.Size;
     MapperArgTypes[I] = C.Type;
     MapperArgNames[I] = C.Name;
   }
 
   int rc = target_data_function(loc, Device, MapperComponents.Components.size(),