[OpenMP5.0] map item can be non-contiguous for target update

[cchen]

Summary

In order not to modify the tgt_target_data_update information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload arg when
the maptype is set as OMP_MAP_DESCRIPTOR. The origin arg is for
passing the pointer information, however, the overloaded arg is an
array of descriptor_dim:

struct descriptor_dim {
  int64_t offset;
  int64_t count;
  int64_t stride
};

and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
arg_size parameter by using dimension size.

More details can be found here: https://github.com/chichunchen/openmp-50-design/blob/master/target_update_noncontiguous.pptx

Edit:
The runtime implementation I'm thinking of is to convert the non-contiguous data into several chunks of contiguous.

For example:

int arr[3][3][3];
#pragma omp target update to (arr[1:2][1:2][0:2])

We can visualize the noncontiguous data as below (X is the data we want to transfer, O is the data want don't bother with):

Dim 0 = {Offset: 0, Count: 1, Stride: 4 bytes (int)}
XXO

Dim 1 = {Offset: 1, Count: 2, Stride: 12 bytes (4 * 3 - since Dim 0 has 3 elements)
OOO
XXO
XXO

Dim 2 = {Offset: 1, Count: 2, Stride: 36 bytes (12 * 3 since Dim 1 has 3 elements)
OOO OOO OOO
OOO XXO XXO
OOO XXO XXO

For the visualization, we know that we want to transfer 4 contiguous chunks and the runtime code could be something similar to:

// we expect this loop to transfer 4 contiguous chunks:
// arr[1][1][0:2]
// arr[1][2][0:2]
// arr[2][1][0:2]
// arr[2][2][0:2]
for (int i = Dim[2].offset; i < Dim[2].count; i++) {
  for (int j = Dim[1].offset; j < Dim[1].count; j++) {
    ptr = bast_ptr + Dim[2].stride * i + Dim[1].stride * j + Dim[2].stride * Dim[0].offset;
    size = Dim[0].count * Dim[0].stride;  // we can hoist it I think
    transfer(ptr, size, /*flag or some other stuff...*/);
  }
}

For this design, we can support strides by just adding an extra dimension. For instance:

int arr[5][5][5]
#pragma omp target update to(arr[0:2:2][1:2:1][0:2:2])

Dim 0 = {offset: 0, count: 1, stride: 4 bytes (int) } // the extra dimension for supporting stride
XO

Dim 1 = {offset: 0, count: 3, stride 8 bytes (4 * 2) }
XOXOX

Dim 2 = {offset: 0, count: 2, stride: 40 bytes (8 * 5) }
OOOOO
XOXOX
XOXOX
OOOOO
OOOOO

Dim 3 = {offset: 0, count: 2, stride: 200 bytes (40 * 5) }
...

[cchen]

Rebase

[cchen]

Remove redundant code

[cchen]

ping

[ABataev]

How are you going to pass this non-contiguous data in the runtime? Are you going to map it in a loop or convert this non-contiguous data into the contiguous and map it as a contiguous chunk of data? Your presentation provides interface only interface changes but has nothing about implementation in the runtime.

[cchen]

In D79972#2049838, @ABataev wrote:

How are you going to pass this non-contiguous data in the runtime? Are you going to map it in a loop or convert this non-contiguous data into the contiguous and map it as a contiguous chunk of data? Your presentation provides interface only interface changes but has nothing about implementation in the runtime.

Hi Alexey, thanks for asking. The runtime implementation I'm thinking of is to convert the non-contiguous data into several chunks of contiguous.

For example:

int arr[3][3][3];

#pragma omp target update to (arr[1:2][1:2][0:2])

We can visualize the noncontiguous data as below (X is the data we want to transfer, O is the data want don't bother with):

Dim 0 = {Offset: 0, Count: 1, Stride: 4bytes (int)}
XXO

Dim 1 = {Offset: 1, Count: 2, Stride: 12bytes (4 * 3 - since Dim 0 has 3 elements)
OOO
XXO
XXO

Dim 2 = {Offset: 1, Count: 2, Stride: 36 bytes (12 * 3 since Dim 1 has 3 elements)
OOO
OOO
OOO
\\\\\
OOO
XXO
XXO
\\\\\
OOO
XXO
XXO

For the visualization, we know that we want to transfer 4 contiguous chunks and the runtime code could be something similar to:

// we expect this loop to transfer 4 contiguous chunks:
// arr[1][1][0:2]
// arr[1][2][0:2]
// arr[2][1][0:2]
// arr[2][2][0:2]
for (int i = Dim[2].offset; i < Dim[2].count; i++) {
  for (int j = Dim[1].offset; j < Dim[1].count; j++) {
    ptr = bast_ptr + Dim[2].stride * i + Dim[1].stride * j + Dim[2].stride * Dim[0].offset;
    size = Dim[0].count * Dim[0].stride;  // we can hoist it I think
    transfer(ptr, size, /*flag or some other stuff...*/);
  }
}

[ABataev]

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

clang/include/clang/AST/OpenMPClause.h

5346

Why do you need this bool flag? Seems to me, it is set to true always if OpenMP >= 50 && Directive == OMPD_target_update. Could check it during the codegen rather than introduce this new extra data here?

[cchen]

clang/include/clang/AST/OpenMPClause.h

5346

You're right, I shouldn't add bool here since we only need it in OMPToClause and OMPFromClause. I was adding it since I'm assuming they should have the same type for the inherited TrailingObject.

[cchen]

In D79972#2058516, @ABataev wrote:

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

My intent is to emit possibly non-contiguous runtime calls only if the analysis in Sema set the IsNonContiguous flag to true.

[ABataev]

In D79972#2058555, @cchen wrote:

In D79972#2058516, @ABataev wrote:

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

My intent is to emit possibly non-contiguous runtime calls only if the analysis in Sema set the IsNonContiguous flag to true.

But this analysis only checks for the directive and the version,nothing else.

[cchen]

In D79972#2058608, @ABataev wrote:

In D79972#2058555, @cchen wrote:

In D79972#2058516, @ABataev wrote:

Is my guess correct that for OpenMP >= 50 for target update directive we always emit possibly non-continuous runtime calls?

My intent is to emit possibly non-contiguous runtime calls only if the analysis in Sema set the IsNonContiguous flag to true.

But this analysis only checks for the directive and the version,nothing else.

The context of the checks for the directive and version:

bool NotWhole =
  checkArrayExpressionDoesNotReferToWholeSize(SemaRef, OASE, CurType);
bool NotUnity =
  checkArrayExpressionDoesNotReferToUnitySize(SemaRef, OASE, CurType);

if (AllowWholeSizeArraySection) {
  // Any array section is currently allowed. Allowing a whole size array
  // section implies allowing a unity array section as well.
  //
  // If this array section refers to the whole dimension we can still
  // accept other array sections before this one, except if the base is a
  // pointer. Otherwise, only unitary sections are accepted.
  if (NotWhole || IsPointer)
    AllowWholeSizeArraySection = false;
} else if (DKind == OMPD_target_update &&
           SemaRef.getLangOpts().OpenMP >= 50) {
  IsNonContiguousRef = true;
} else if (AllowUnitySizeArraySection && NotUnity) {
  // A unity or whole array section is not allowed and that is not
  // compatible with the properties of the current array section.
  SemaRef.Diag(
    ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
    << OASE->getSourceRange();
  return false;
}

The original analysis checks for non-contiguous by finding if there is more than one "array-section" expression with length greater than one. Therefore, I added my check there to allow more than one array-section with length greater than one by depending on the existing analysis (and also set IsNonContiguous to true so that we can pass it to codegen rather than doing analysis in codegen). This change allows me to pass all the existing lit test but still emit the "non-contiguous" runtime.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

16621

@ABataev , I guess you're saying the condition should be !AllowWholeSizeArraySection && DKind == OMPD_target_update && SemaRef.getLangOpts().OpenMP >= 50?

[ABataev]

Did you think about implementing it in the compiler instead of the runtime?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7624–7627

Do you really need to count DimSize for array shaping operators and array subscript expressions? I don't see tests for it.

clang/lib/Sema/SemaOpenMP.cpp

16621

No, what I want is to try to simplify the code. I see now why do you need this flag. I'm just thinking can we avoid adding this flag to the clause and save some mem space?

[cchen]

Fix based on feedback

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7624–7627

You're right, I don't need to count DimSize for array shaping and array subscript.

clang/lib/Sema/SemaOpenMP.cpp

16621

But we also don't want to do the analysis in codegen I guess? Also, if we emit non-contiguous runtime for every target update call, we need to change tons of stuff (tons of lit tests, runtime implementation, etc...).

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

16621

Maybe make it a part of MappableComponent, if possible, and put it into PointerIntPair<Expr *, 1, bool> AssociatedExpression;?

[cchen]

• Use PointerIntPair to pass non-contiguous information in AST
• Error out in Sema if we don't have enough size information for cases involving pointers
• Allows *arr[N][M] since we don't need size information for the last dimension
• Add more test cases

[ABataev]

Still: Did you think about implementing it in the compiler instead of the runtime?

clang/include/clang/AST/OpenMPClause.h

4752–4753

I would suggest to pass Expr * and bool as separate parameters here rather than as PointerIntPair

4766

isNonContiguous()

clang/lib/Sema/SemaOpenMP.cpp

16442

Add default initializer

16623–16627

Remove braces here, they are not needed.

clang/lib/Serialization/ASTWriter.cpp

6584

There is a member function writeBool

6609

Same, use writeBool

[cchen]

In D79972#2068976, @ABataev wrote:

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

[ABataev]

In D79972#2069322, @cchen wrote:

In D79972#2068976, @ABataev wrote:

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

[cchen]

In D79972#2069358, @ABataev wrote:

In D79972#2069322, @cchen wrote:

In D79972#2068976, @ABataev wrote:

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

Oh, I would prefer to call tgt_target_data_update once in the compiler and I'm also doing it now.

[ABataev]

In D79972#2069366, @cchen wrote:

In D79972#2069358, @ABataev wrote:

In D79972#2069322, @cchen wrote:

In D79972#2068976, @ABataev wrote:

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

Oh, I would prefer to call tgt_target_data_update once in the compiler and I'm also doing it now.

I was not quite correct. What I mean, is to generate the array with the array section as VLA in the compiler, and fill it in the loop generated by the compiler for non-contiguous sections but not in the runtime?
Say, we have the code:

int arr[3][3]
...
 #pragma omp update to(arr[1:2][1:2]

In this case, we're going to transfer the next elements:

000
0xx
0xx

In the compiler-generated code we emit something like this:

void *bptr[<n>];
void *ptr[<n>];
int64 sizes[<n>];
int64 maptypes[<n>];
for (int i = 0; i < <n>; ++i) {
  bptr[i] = &arr[1+i][1];
  ptr[i] = &arr[1+i][1];
  sizes[i] = ...;'
  maptypes[i] = ...;
}
call void @__tgt_target_data_update(i64 -1, i32 <n>, bptr, ptr, sizes, maptypes);

With this solution, you won't need to modify the runtime and add a new mapping flag.

[cchen]

Fix based on feedback

[cchen]

In D79972#2069435, @ABataev wrote:

In D79972#2069366, @cchen wrote:

In D79972#2069358, @ABataev wrote:

In D79972#2069322, @cchen wrote:

In D79972#2068976, @ABataev wrote:

Still: Did you think about implementing it in the compiler instead of the runtime?

I'm not sure I understand your question, which part of code are you asking?
The main work compiler needs to do is to send the {offset, count, stride} struct to runtime.

I mean did you think about calling __tgt_target_data_update function in a loop in the compiler-generated code instead of putting it into the runtime?

Oh, I would prefer to call tgt_target_data_update once in the compiler and I'm also doing it now.

I was not quite correct. What I mean, is to generate the array with the array section as VLA in the compiler, and fill it in the loop generated by the compiler for non-contiguous sections but not in the runtime?
Say, we have the code:

int arr[3][3]
...
 #pragma omp update to(arr[1:2][1:2]

In this case, we're going to transfer the next elements:

000
0xx
0xx

In the compiler-generated code we emit something like this:

void *bptr[<n>];
void *ptr[<n>];
int64 sizes[<n>];
int64 maptypes[<n>];
for (int i = 0; i < <n>; ++i) {
  bptr[i] = &arr[1+i][1];
  ptr[i] = &arr[1+i][1];
  sizes[i] = ...;'
  maptypes[i] = ...;
}
call void @__tgt_target_data_update(i64 -1, i32 <n>, bptr, ptr, sizes, maptypes);

With this solution, you won't need to modify the runtime and add a new mapping flag.

For my current implementation, we have discussed in the bi-weekly meeting several weeks back, and there was a general consensus that it was an acceptable approach.

The major advantage of sending a descriptor to runtime can be elaborated in the following example:

#define N 10000
int a[N][2];
…
#pragma amp target update to (a[0:N][0:1])

This would require passing through O(N) entries in the tgt_target_data_update call, or 10000 entries. The current implementation only require a descriptor with 2 entries. I think this could be a real concern -
splitting out the transfers in compiler-generated code results in a list containing one entry per non-contiguous chunk (easily hitting scaling issues), while the descriptor approach is bounded by the number of dimensions.
That seems like a pretty compelling reason to use the descriptor - it’s much more space efficient.

Also, the descriptor idea is very similar to how Cray supported Fortran dope vectors for years (we send in a pointer to a dope vector rather than a pointer to the data, and a flag to indicate it’s a dope vector, and the runtime library handles it as a dope vector).
I think the runtime library changes will not be very extensive or difficult at all and we’re very willing to implement the runtime for non-contiguous.

[cchen]

ping

[ABataev]

Do you have a test for mapping of something like arr[0][:n], where the base is an array subscript and the remaining part is an array section?

clang/include/clang/AST/OpenMPClause.h

4756–4757

I think you can initialize AssociatedExpressionNonContiguousPr using just AssociatedExpressionNonContiguousPr(AssociatedExpression, IsNonContiguous) form, no?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7114

Restore original formatting

7602–7603

Better to convert it to !IsNonContiguous && isFinalArraySectionExpression(I->getAssociatedExpression()).

7622

Use prefix form ++DimSize.

7679

No need for parameter name comment here, it is required only if the true|false constants are used

7736

Same, comment not required

7913–7927

Can we merge the functionality in this new function with the existing ones somehow? It is not the best idea to duplicate functionality using copy-paste if any.

8583–8585

Why removed the comment?

8950

Same question as before - can we merge this functionality with the existing functions?

clang/lib/Sema/SemaOpenMP.cpp

18468–18473

Use .emplace_back(SimpleRefExpr, D, false);

clang/lib/Serialization/ASTReader.cpp

12514–12516

.emplace_back(AssociatedExprPr, AssociatedDecl, /*IsNonContiguous=*/false);

12632–12634

Same, use emplace_back()

12683–12684

Same, use emplace_back()

12733–12734

Same, use emplace_back()

12819–12820

Same, use emplace_back()

clang/test/OpenMP/target_update_to_messages.cpp

147

Delete this extra line

[cchen]

Fix based on feedback

[cchen]

In D79972#2082017, @ABataev wrote:

Do you have a test for mapping of something like arr[0][:n], where the base is an array subscript and the remaining part is an array section?

I'm not having it right now, but it seems like if the base is an array subscript and the remaining part is an array section, then this map-item will always be contiguous, and will not trigger my code in Codegen. I can still add a test for Sema though.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7049

I'm not sure about the value of this flag. If I do recall it correctly, this value might be used for something different by XL compiler, for example. Maybe use some other value, maybe high bits? It is a kind of service flag, not data mapping attribute, so better to move it to high bits (the bit before OMP_MAP_MEMBER_OF maybe?).

7806–7807

Use range-based loop, if possible.

7809–7811

Do you really need to analyze array subscript expressions here? I though that we should analyze only array sections, no?

7868

Same, try to use range-based loop, if possible.

7873

Same question about array subscript expressions.

8218–8223

Just `generateInfoForComponentList(

L.MapType, L.MapModifiers, L.Components, CurBasePointers,
CurPointers, CurSizes, CurTypes, CurDims, PartialStruct,
IsFirstComponentList, L.IsImplicit,
/*OverlappedElements=*/llvm::None,
L.Components.back().isNonContiguous(), &CurOffsets, &CurCounts, &CurStrides);`

8768–8771

Can we encapsulate these new data into CGOpenMPRuntime::TargetDataInfo?

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7327–7330

I would prefer to pack these 4 params into a single parameter (a struct). Also, can we put Dims parameter into the list of the optional parameters?

7805

Expand auto here to a real type

7821

What if the base is a pointer, not an array?

7831–7838

The code for SizeV must be under the control of the next if:

if (DimSizes.size() < Components.size() - 1) {
 ....
}

7834

Create directly as of CGF.Int64Ty type.

7859

Expand auto here to a real type

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7807

Can we have anything else except for array section here? If not, use just cast. If yes, use continue to simplify complexity:

if (!OASE)
  continue;
...

7861–7864

Can we have anything else except for array section here? If not, use just cast. If yes, use continue to simplify complexity:

if (!OASE)
  continue;
...

7872–7873

Do you really to pass real offsets here? Can we use pointers instead?

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7821

The if (ElementType) condition only push back stride when base is not pointer. I'm now allowing one dimension size to be unknown (pointer as base) and sema has analysis to check if more than one indirection as base. My last codegen test case is for testing pointer as base.

7831–7838

I don't think I understand this one. Why do you remove SizeV in the if condition?

7834

Doing this I'll get assertion error in this exact line if on a 32-bits target.

7872–7873

Do you mean I should set the type of Offset to Expr*?

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7327

Can we encapsulate Dims into StructNonContiguousInfo?

7831–7832

If only CAT or VAT is allowed, then transform this if into:

else {
  assert(VAT&& ...);

7831–7838

This is for SizeV. You don't use it if DimSizes.size() < Components.size() - 1 is false, looks like memory leak.

7834

Hmm, why, can you investigate?

7869–7872

Can we have anything else except for array section here? If not, use just cast. If yes, use continue to simplify complexity:

if (!OASE)
  continue;
...

7872–7873

Currently, you're passing offsets to the runtime. Can we pass pointers instead? I mean, for a[b] you pass b to the runtime, can we pass &a[b] instead?

7911

Avoid expressions with some side effects, like *DI++

[cchen]

Resolve issues

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7834

My comment was not accurate, I've updated it. What I want to convey is that we can only have CAT, VAT, or pointer here, since analysis in Sema has a restriction for it. (SemaOpenMP line 16623)

7869–7872

Not sure about this one, I've added:

if (!OASE)
  continue;
...

7872–7873

Yes, I'm fine either passing index or passing address, though I'm curious why you're recommending passing address.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7834

It does not relate to the comments thread but I got it. Anyway, try to investigate why the compiler crashes if you try to cr4eate a constant ща ]СПАюШте64Ен] directly.

7872–7873

It is going to simplify the codegen. Currently, to get the offset, you need to dig through all the elements of the array section. If, instead, you use the pointers, you would not need to do this and you can rely on something like CGF.EmitArraySectionLValue(). At least, I hope so.

7903

The check is not required, you already checked that the expression must be array section only.

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7834

I'll investigate it, thanks.

7872–7873

After discussed with my colleagues, I think passing relative offset makes more sense.

For a 1-dim array, storing the offset as a pointer could work, but it seems strange to me to store as a pointer when there are 2+ dimensions with multiple disjoint chunks of memory because the pointer can only point to the offset for the first chunk. That is, a pointer would refer to an absolute location in a single chunk, whereas the offset is relative to the start of any chunk.

For example:

int a[4][4];
#pragma omp target update to(a[1:2][1:2])

This is two disjoint chunks of memory:

XXXX
XOOX
XOOX
XXXX

The offset for the outer dimension could be store as a pointer, since there is only one instance of that dimension:

Dim1: Offset=&a[1]

But, the inner dimension is "instantiated" twice, once for each element in the outer dimension. So, there are really two absolute pointers, depending on which instance (element in the outer dimension) you're talking about:

Dim2: Offset=&a[1][1]
Dim2: Offset=&a[2][1]

We could set the policy that the absolute offset would always be expressed as the offset in the first instance, but then wouldn't we need to refer to that location when computing the offset for all of the other instances? That seems unintuitive to me, and potentially complicates the implementation. The relative offset makes a lot more senes to me - for a starting point, what relative offset is needed for each dimension. The starting point for the outermost dimension does require the base address, but all inner dimensions have a variable starting pointer based on which element in the outer dimensions you're currently looking at.

[cchen]

Fix Int64Ty issue (The bitNum of APInt I used before is 32)

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7920

Use preincrement

8224

No need to add /*OverlappedElements=*/llvm::None here, it is default value.

8878–8880

Better just to have something like this:

if (!IsNonContiguous || Info.Offsets.empty() || Info.NumberOfPtrs == 0)
  return;
...

[cchen]

Fix based on feedback

[RaviNarayanaswamy]

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

[cchen]

In D79972#2104854, @RaviNarayanaswamy wrote:

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

There's only one runtime call for your case. and args will be { descriptor_1, x, descriptor_2 }, where descriptor_1 will be { { 1, 2, 80 }, { 1, 2, 20 }, { 0, 2, 4 } }, descriptor_2 will be { { 1, 5, 16 }, { 0, 2, 4 } }. There's analysis in Sema that detecting if the item is non-contiguous or not and codegen only generate descriptor for non-contiguous item.

[cchen]

Updated test for clarification

[cchen]

In D79972#2104854, @RaviNarayanaswamy wrote:

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

I have added a test basically base on the case in your comment (CK19 in target_update_codegen.cpp). Thanks.

[cchen]

ping

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7842–7844

No need for braces here

8935

C.getTypeAlignInChars(C.VoidPtrTy)->CGF.getPointerAlign()

10317–10320

Better just to pass Info.Offsets, Info.Counts and Info.Strides as arguments to generateAllInfo() function and do not create local copies at all.

10570–10573

Same, pass the fields as arguments instead.

clang/lib/Serialization/ASTReader.cpp

12515–12516

Still calling an extra constructor here, just .emplace_back(AssociatedExprPr, AssociatedDecl, /*IsNonContiguous=*/false);

12633–12635

Just Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);

12684–12685

.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);

12734–12735

.emplace_back(AssociatedExprPr, AssociatedDecl, /*IsNonContiguous=*/false);

12776–12778

.emplace_back(AssociatedExpr, AssociatedDecl, /*IsNonContiguous*/ false);

12819–12821

.emplace_back(AssociatedExpr, AssociatedDecl, /*IsNonContiguous=*/false));

[cchen]

Fix coding style

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8761–8763

Do you really need to pass Dims here if you have Dims data member in Info parameter? Why you can't use Info.Dims instead?

8881–8909

Maybe worth it to outline it into a separate function to reduce code size and the complexity of this function? And just call this new function here.

clang/lib/Sema/SemaOpenMP.cpp

16624

Better to use integer value as selectors, not boolean.

18519–18523

.emplace_back(SimpleRefExpr, D, /*IsNonContiguous=*/false);

18588

Add a comment for false argument with the name of parameter.

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8761–8763

I think I haven't added Dims in TargetDataInfo atm, I'll add into it and then use it via Info.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

16624

The selector for err_omp_section_length_undefined is a bool value. (true for unknown bound false for not a array type, so always be true here).
Do you mean that I need to create a new kind of diagnosis message here and use integer as selectors?

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

16624

No, it is an integer, starts from 0

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10315–10316

Better to pass Info here directly.

[cchen]

Pass Info directly

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

16664

/*IsNonContiguous=*/false

16684

/*IsNonContiguous=*/false

18469–18473

.emplace_back(SimpleRefExpr, D, /*IsNonContiguous=*/false);

[cchen]

Fix based on feedback

[cchen]

Rebase and resolve conflictions

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7049

Hi @ABataev, is there any place I can find which value has been used for lower bits (like 0x800, 0x1000)?

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7049

I rather doubt. You can try to ask @kkwli0

[kkwli0]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7049

We are using 0x800. I think your current choice should be fine.

[cchen]

@ABataev , I'm considering emitting an extra dimension for a non-contiguous descriptor to support stride in this patch (stride = 1 in array section is just a special case for computing stride, however, the formula computing stride do not change). Do you think I should do it in this patch?

Computing of stride after support stride in array section:

int arr[5][5][5];
#pragma omp target update to(arr[0:2:2][1:2:1][0:2:2]

D0: { offset = 0, count = 1, stride = 4 }                                           // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 4 * 1 * 2 = 8 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }             // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 2 = 200 }      // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200

For the case in this patch (stride = 1), we can use the same formula for computing stride with extra dimension:

int arr[5][5][5];
#pragma omp target update to(arr[0:2][1:2][0:2]

D0: { offset = 0, count = 1, stride = 4 }                                          // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 4 * 1 * 1 = 4 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 1 = 4
D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }            // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 1 = 100 }     // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 1 = 100

The extra dimension does not affect the runtime implementation at all since runtime will try to merge inner dimensions if they are contiguous. Take the above case for example (arr[0:2][1:2][0:2]):
The product of count and stride for D0 is 4 which is the same as the stride of D1, therefore, runtime just ignores D0.

[ABataev]

In D79972#2124108, @cchen wrote:

@ABataev , I'm considering emitting an extra dimension for a non-contiguous descriptor to support stride in this patch (stride = 1 in array section is just a special case for computing stride, however, the formula computing stride do not change). Do you think I should do it in this patch?

Computing of stride after support stride in array section:

int arr[5][5][5];
#pragma omp target update to(arr[0:2:2][1:2:1][0:2:2]

D0: { offset = 0, count = 1, stride = 4 }                                           // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 4 * 1 * 2 = 8 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }             // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 2 = 200 }      // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200

For the case in this patch (stride = 1), we can use the same formula for computing stride with extra dimension:

int arr[5][5][5];
#pragma omp target update to(arr[0:2][1:2][0:2]

D0: { offset = 0, count = 1, stride = 4 }                                          // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 4 * 1 * 1 = 4 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 1 = 4
D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }            // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 1 = 100 }     // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 1 = 100

The extra dimension does not affect the runtime implementation at all since runtime will try to merge inner dimensions if they are contiguous. Take the above case for example (arr[0:2][1:2][0:2]):
The product of count and stride for D0 is 4 which is the same as the stride of D1, therefore, runtime just ignores D0.

You can do this patch. But at first, you need to commit the runtime part of the patch that supports it, and the part that introduces stride support.

[cchen]

Created a new patch with the support for stride: https://reviews.llvm.org/D84192.

clang/include/clang/AST/OpenMPClause.h

 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtIterator.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Basic/SourceLocation.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/MapVector.h"
+#include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/iterator.h"
 #include "llvm/ADT/iterator_range.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include "llvm/Frontend/OpenMP/OMPContext.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/Support/TrailingObjects.h"

 public:
   /// Class that represents a component of a mappable expression. E.g.
   /// for an expression S.a, the first component is a declaration reference
   /// expression associated with 'S' and the second is a member expression
   /// associated with the field declaration 'a'. If the expression is an array
   /// subscript it may not have any associated declaration. In that case the
   /// associated declaration is set to nullptr.
   class MappableComponent {
-    /// Expression associated with the component.
-    Expr *AssociatedExpression = nullptr;
+    /// Pair of Expression and Non-contiguous pair  associated with the
+    /// component.
+    llvm::PointerIntPair<Expr *, 1, bool> AssociatedExpressionNonContiguousPr;
 
     /// Declaration associated with the declaration. If the component does
     /// not have a declaration (e.g. array subscripts or section), this is set
     /// to nullptr.
     ValueDecl *AssociatedDeclaration = nullptr;
 
   public:
     explicit MappableComponent() = default;
     explicit MappableComponent(Expr *AssociatedExpression,
-                               ValueDecl *AssociatedDeclaration)
-        : AssociatedExpression(AssociatedExpression),
+                               ValueDecl *AssociatedDeclaration,
+                               bool IsNonContiguous)
+        : AssociatedExpressionNonContiguousPr(AssociatedExpression,
+                                              IsNonContiguous),
           AssociatedDeclaration(
               AssociatedDeclaration
                   ? cast<ValueDecl>(AssociatedDeclaration->getCanonicalDecl())
                   : nullptr) {}
 
-    Expr *getAssociatedExpression() const { return AssociatedExpression; }
+    Expr *getAssociatedExpression() const {
+      return AssociatedExpressionNonContiguousPr.getPointer();
+    }
+
+    bool isNonContiguous() const {
+      return AssociatedExpressionNonContiguousPr.getInt();
+    }
 
     ValueDecl *getAssociatedDeclaration() const {
       return AssociatedDeclaration;
     }
   };
 
   // List of components of an expression. This first one is the whole
   // expression and the last one is the base expression.

 /// In this example directive '#pragma omp target' has clause 'map'
 /// with the variables 'a' and 'b'.
 class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
                            private llvm::TrailingObjects<
                                OMPMapClause, Expr *, ValueDecl *, unsigned,
                                OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// Define the sizes of each trailing object array except the last one. This
   /// is required for TrailingObjects to work properly.
   size_t numTrailingObjects(OverloadToken<Expr *>) const {
     // There are varlist_size() of expressions, and varlist_size() of
     // user-defined mappers.
     return 2 * varlist_size();

clang/lib/CodeGen/CGOpenMPRuntime.h

public:
     llvm::Value *SizesArray = nullptr;
     /// The array of map types passed to the runtime library.
     llvm::Value *MapTypesArray = nullptr;
     /// The total number of pointers passed to the runtime library.
     unsigned NumberOfPtrs = 0u;
     /// Map between the a declaration of a capture and the corresponding base
     /// pointer address where the runtime returns the device pointers.
     llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
+    /// The array of array of offsets passed to the runtime library.
+    SmallVector<SmallVector<llvm::Value *, 4>, 4> Offsets;
+    /// The array of array of counts passed to the runtime library.
+    SmallVector<SmallVector<llvm::Value *, 4>, 4> Counts;
+    /// The array of array of strides passed to the runtime library.
+    SmallVector<SmallVector<llvm::Value *, 4>, 4> Strides;
 
     explicit TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo)
         : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}
     /// Clear information about the data arrays.
     void clearArrayInfo() {
       BasePointersArray = nullptr;
       PointersArray = nullptr;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

enum OpenMPOffloadMappingFlags : uint64_t {
     OMP_MAP_PRIVATE = 0x80,
     /// Pass the element to the device by value.
     OMP_MAP_LITERAL = 0x100,
     /// Implicit map
     OMP_MAP_IMPLICIT = 0x200,
     /// Close is a hint to the runtime to allocate memory close to
     /// the target device.
     OMP_MAP_CLOSE = 0x400,
+    /// Signal that the runtime library should use args as an array of
+    /// descriptor_dim pointers and use args_size as dims. Used when we have
+    /// non-contiguous list items in target update directive
+    OMP_MAP_DESCRIPTOR = 0x100000000000,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
     LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */ OMP_MAP_MEMBER_OF),
   };
 
   /// Get the offset of the OMP_MAP_MEMBER_OF field.
   static unsigned getFlagMemberOffset() {
public:
     llvm::Value *operator*() const { return Ptr; }
     const ValueDecl *getDevicePtrDecl() const { return DevPtrDecl; }
     void setDevicePtrDecl(const ValueDecl *D) { DevPtrDecl = D; }
   };
 
   using MapBaseValuesArrayTy = SmallVector<BasePointerInfo, 4>;
   using MapValuesArrayTy = SmallVector<llvm::Value *, 4>;
   using MapFlagsArrayTy = SmallVector<OpenMPOffloadMappingFlags, 4>;
+  using MapDimArrayTy = SmallVector<uint64_t, 4>;
+  using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
 
   /// Map between a struct and the its lowest & highest elements which have been
   /// mapped.
   /// [ValueDecl *] --> {LE(FieldIndex, Pointer),
   ///                    HE(FieldIndex, Pointer)}
   struct StructRangeInfoTy {
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> LowestElem = {
         0, Address::invalid()};
struct MapInfo {
     bool ReturnDevicePointer = false;
     bool IsImplicit = false;
 
     MapInfo() = default;
     MapInfo(
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
         OpenMPMapClauseKind MapType,
         ArrayRef<OpenMPMapModifierKind> MapModifiers,
         bool ReturnDevicePointer, bool IsImplicit)
         : Components(Components), MapType(MapType), MapModifiers(MapModifiers),
           ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
   };
 
   /// If use_device_ptr is used on a pointer which is a struct member and there
   /// is no map information about it, then emission of that entry is deferred
   /// until the whole struct has been processed.
   struct DeferredDevicePtrEntryTy {
llvm::Value *getExprTypeSize(const Expr *E) const {
     }
     return CGF.getTypeSize(ExprTy);
   }
 
   /// Return the corresponding bits for a given map clause modifier. Add
   /// a flag marking the map as a pointer if requested. Add a flag marking the
   /// map as the first one of a series of maps that relate to the same map
   /// expression.
-  OpenMPOffloadMappingFlags getMapTypeBits(
-      OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
-      bool IsImplicit, bool AddPtrFlag, bool AddIsTargetParamFlag) const {
+  OpenMPOffloadMappingFlags
+  getMapTypeBits(OpenMPMapClauseKind MapType,
+                 ArrayRef<OpenMPMapModifierKind> MapModifiers, bool IsImplicit,
+                 bool AddPtrFlag, bool AddIsTargetParamFlag,
+                 bool IsNonContiguous) const {
     OpenMPOffloadMappingFlags Bits =
         IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE;
     switch (MapType) {
     case OMPC_MAP_alloc:
     case OMPC_MAP_release:
       // alloc and release is the default behavior in the runtime library,  i.e.
       // if we don't pass any bits alloc/release that is what the runtime is
       // going to do. Therefore, we don't need to signal anything for these two
getMapTypeBits(OpenMPMapClauseKind MapType,
     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 (IsNonContiguous)
+      Bits |= OMP_MAP_DESCRIPTOR;
     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.
   bool isFinalArraySectionExpression(const Expr *E) const {
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
 
bool isFinalArraySectionExpression(const Expr *E) const {
     return ConstLength.getSExtValue() != 1;
   }
 
   /// Generate the base pointers, section pointers, sizes and map type
   /// bits for the provided map type, map modifier, and expression components.
   /// \a IsFirstComponent should be set to true if the provided set of
   /// components is the first associated with a capture.
   void generateInfoForComponentList(
-      OpenMPMapClauseKind MapType,
-      ArrayRef<OpenMPMapModifierKind> MapModifiers,
+      OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
       MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
-      MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+      MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, MapDimArrayTy &Dims,
       StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
       bool IsImplicit,
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
-          OverlappedElements = llvm::None) const {
+          OverlappedElements = llvm::None,
+      bool IsNonContiguous = false,
+      MapNonContiguousArrayTy *const Offsets = nullptr,
+      MapNonContiguousArrayTy *const Counts = nullptr,
+      MapNonContiguousArrayTy *const Strides = nullptr) const {
     // The following summarizes what has to be generated for each map and the
     // types below. The generated information is expressed in this order:
     // base pointer, section pointer, size, flags
     // (to add to the ones that come from the map type and modifier).
     //
     // double d;
     // int i[100];
     // float *p;
void generateInfoForComponentList(
     // in the component list which is a member expression. Useful when we have a
     // pointer or a final array section, in which case it is the previous
     // component in the list which tells us whether we have a member expression.
     // E.g. X.f[:]
     // While processing the final array section "[:]" it is "f" which tells us
     // whether we are dealing with a member of a declared struct.
     const MemberExpr *EncounteredME = nullptr;
 
+    // Track for the total number of dimension.
+    uint64_t DimSize = 0;
+
     for (; I != CE; ++I) {
       // If the current component is member of a struct (parent struct) mark it.
       if (!EncounteredME) {
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
         if (EncounteredME)
           ShouldBeMemberOf = true;
       }
 
       auto Next = std::next(I);
 
       // We need to generate the addresses and sizes if this is the last
       // component, if the component is a pointer or if it is an array section
       // whose length can't be proved to be one. If this is a pointer, it
       // becomes the base address for the following components.
 
       // A final array section, is one whose length can't be proved to be one.
+      // If the map item is non-contiguous then we don't treat any array section
+      // as final array section.
       bool IsFinalArraySection =
+          !IsNonContiguous &&
           isFinalArraySectionExpression(I->getAssociatedExpression());
 
       // Get information on whether the element is a pointer. Have to do a
       // special treatment for array sections given that they are built-in
       // types.
       const auto *OASE =
           dyn_cast<OMPArraySectionExpr>(I->getAssociatedExpression());
       const auto *OAShE =
           dyn_cast<OMPArrayShapingExpr>(I->getAssociatedExpression());
       const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
       const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
       bool IsPointer =
           OAShE ||
           (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
                        .getCanonicalType()
                        ->isAnyPointerType()) ||
           I->getAssociatedExpression()->getType()->isAnyPointerType();
-      bool IsNonDerefPointer = IsPointer && !UO && !BO;
+      bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
+
+      if (OASE)
+        ++DimSize;
 
       if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
         // If this is not the last component, we expect the pointer to be
         // associated with an array expression or member expression.
         assert((Next == CE ||
                 isa<MemberExpr>(Next->getAssociatedExpression()) ||
                 isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
                 isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
                 isa<UnaryOperator>(Next->getAssociatedExpression()) ||
                 isa<BinaryOperator>(Next->getAssociatedExpression())) &&
                "Unexpected expression");
 
         Address LB = Address::invalid();
for (; I != CE; ++I) {
                   PartialStruct.HighestElem.first)>::max(),
               HB};
           PartialStruct.Base = BP;
           // Emit data for non-overlapped data.
           OpenMPOffloadMappingFlags Flags =
               OMP_MAP_MEMBER_OF |
               getMapTypeBits(MapType, MapModifiers, IsImplicit,
                              /*AddPtrFlag=*/false,
-                             /*AddIsTargetParamFlag=*/false);
+                             /*AddIsTargetParamFlag=*/false, IsNonContiguous);
           LB = BP;
           llvm::Value *Size = nullptr;
           // Do bitcopy of all non-overlapped structure elements.
           for (OMPClauseMappableExprCommon::MappableExprComponentListRef
                    Component : OverlappedElements) {
             Address ComponentLB = Address::invalid();
             for (const OMPClauseMappableExprCommon::MappableComponent &MC :
                  Component) {
               if (MC.getAssociatedDeclaration()) {
                 ComponentLB =
                     CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
                         .getAddress(CGF);
                 Size = CGF.Builder.CreatePtrDiff(
                     CGF.EmitCastToVoidPtr(ComponentLB.getPointer()),
                     CGF.EmitCastToVoidPtr(LB.getPointer()));
                 break;
               }
             }
             BasePointers.push_back(BP.getPointer());
             Pointers.push_back(LB.getPointer());
             Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
                                                       /*isSigned=*/true));
             Types.push_back(Flags);
+            Dims.push_back(IsNonContiguous ? DimSize : 0);
             LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
           }
           BasePointers.push_back(BP.getPointer());
           Pointers.push_back(LB.getPointer());
           Size = CGF.Builder.CreatePtrDiff(
               CGF.EmitCastToVoidPtr(
                   CGF.Builder.CreateConstGEP(HB, 1).getPointer()),
               CGF.EmitCastToVoidPtr(LB.getPointer()));
           Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           Types.push_back(Flags);
+          Dims.push_back(IsNonContiguous ? DimSize : 0);
           break;
         }
         llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
         if (!IsMemberPointer) {
           BasePointers.push_back(BP.getPointer());
           Pointers.push_back(LB.getPointer());
           Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+          Dims.push_back(IsNonContiguous ? DimSize : 0);
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags = getMapTypeBits(
               MapType, MapModifiers, IsImplicit,
               !IsExpressionFirstInfo || RequiresReference,
-              IsCaptureFirstInfo && !RequiresReference);
+              IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
 
           if (!IsExpressionFirstInfo) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
             // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
             if (IsPointer)
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
                          OMP_MAP_DELETE | OMP_MAP_CLOSE);
 
             if (ShouldBeMemberOf) {
               // Set placeholder value MEMBER_OF=FFFF to indicate that the flag
               // should be later updated with the correct value of MEMBER_OF.
               Flags |= OMP_MAP_MEMBER_OF;
for (; I != CE; ++I) {
         // The pointer becomes the base for the next element.
         if (Next != CE)
           BP = LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
       }
     }
+
+    if (IsNonContiguous) {
+      const ASTContext &Context = CGF.getContext();
+
+      MapValuesArrayTy CurOffsets;
+      MapValuesArrayTy CurCounts;
+      MapValuesArrayTy CurStrides;
+      llvm::Value *CurStride = nullptr;
+      SmallVector<llvm::Value *, 4> DimSizes;
+
+      // Collect Size information for each dimension and get the element size as
+      // the first Stride. For example, for `int arr[10][10]`, the DimSizes
+      // should be [10, 10] and the first stride is 4 btyes.
+      for (const auto &Component : Components) {
+        const Expr *AssocExpr = Component.getAssociatedExpression();
+        const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+        if (OASE) {
+          QualType Ty;
+          Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+          auto *CAT = Context.getAsConstantArrayType(Ty);
+          auto *VAT = Context.getAsVariableArrayType(Ty);
+          // Get element size if CurStrides is empty.
+          if (CurStrides.empty()) {
+            const Type *ElementType = nullptr;
+            uint64_t ElementTypeSize;
+            if (CAT) {
+              ElementType = CAT->getElementType().getTypePtr();
+              ElementTypeSize =
+                  Context.getTypeSizeInChars(ElementType).getQuantity();
+            } else if (VAT) {
+              ElementType = VAT->getElementType().getTypePtr();
+              ElementTypeSize =
+                  Context.getTypeSizeInChars(ElementType).getQuantity();
+            }
+            if (ElementType)
+              CurStrides.push_back(
+                  llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize));
+          }
+          // Get dimension value.
+          llvm::Value *SizeV = nullptr;
+          if (CAT) {
+            llvm::APInt Size = CAT->getSize();
+            SizeV = llvm::ConstantInt::get(CGF.SizeTy, Size);
+          } else if (VAT) {
+            const Expr *Size = VAT->getSizeExpr();
+            SizeV = CGF.EmitScalarExpr(Size);
+          }
+          // We need all the dimension size except for the last dimension.
+          assert((VAT || CAT || &Component == &*Components.begin()) &&
+                 "Should be either ConstantArray or VariableArray if not the "
+                 "first Component");
+          if (SizeV && DimSizes.size() < Components.size() - 1)
+            DimSizes.push_back(CGF.Builder.CreateIntCast(SizeV, CGF.Int64Ty,
+                                                         /*IsSigned=*/false));
+        }
+      }
+
+      // We need dimension size to compute stride
+      auto DI = DimSizes.begin();
+
+      // Collect info for non-contiguous. Notice that offset, count, and stride
+      // are only meaningful for array-section, so we insert a null for anything
+      // other than array-section.
+      // Also, the size of offset, count, and stride are not the same as
+      // pointers, base_pointers, sizes, or dims. Instead, the size of offset,
+      // count, and stride are the same as the number of non-contiguous
+      // declaration in target update to/from clause.
+      for (const auto &Component : Components) {
+        const Expr *AssocExpr = Component.getAssociatedExpression();
+        const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+        if (OASE) {
+          // Offset
+          const Expr *OffsetExpr = nullptr;
+          OffsetExpr = OASE->getLowerBound();
+          llvm::Value *Offset = nullptr;
+          if (!OffsetExpr) {
+            // If offset is absent, then we just set it to zero.
+            Offset = llvm::ConstantInt::get(CGF.Int64Ty, 0);
+          } else {
+            Offset = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(OffsetExpr),
+                                               CGF.Int64Ty,
+                                               /*isSigned=*/false);
+          }
+          CurOffsets.push_back(Offset);
+
+          // Count
+          const Expr *CountExpr = nullptr;
+          if (OASE)
+            CountExpr = OASE->getLength();
+          llvm::Value *Count = nullptr;
+          if (!CountExpr) {
+            // If length is absent then we calculate it as (Total length -
+            // lower_bound)
+            Count = CGF.Builder.CreateNUWSub(*DI, Offset);
+          } else {
+            Count = CGF.EmitScalarExpr(CountExpr);
+          }
+          Count =
+              CGF.Builder.CreateIntCast(Count, CGF.Int64Ty, /*isSigned=*/false);
+          CurCounts.push_back(Count);
+
+          // Stride = previous stride * previous dimension size
+          // Take `int arr[5][10]` and `arr[0:2][0:2]` as an example:
+          //              Dimension 1       Dimension 0
+          //    Offset    0                 0
+          //    Count     2                 2
+          //    Stride    40 bytes (4x10)   4 bytes (int)
+          if (DI != DimSizes.end()) {
+            CurStride = CGF.Builder.CreateNUWMul(CurStrides.back(), *DI++);
+            CurStrides.push_back(CurStride);
+          }
+        }
+      }
+
+      Offsets->push_back(CurOffsets);
+      Counts->push_back(CurCounts);
+      Strides->push_back(CurStrides);
+    }
   }
 
   /// Return the adjusted map modifiers if the declaration a capture refers to
   /// appears in a first-private clause. This is expected to be used only with
   /// directives that start with 'target'.
   MappableExprsHandler::OpenMPOffloadMappingFlags
   getMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap) const {
     assert(Cap.capturesVariable() && "Expected capture by reference only!");
 
     // A first private variable captured by reference will use only the
     // 'private ptr' and 'map to' flag. Return the right flags if the captured
     // declaration is known as first-private in this handler.
     if (FirstPrivateDecls.count(Cap.getCapturedVar())) {
       if (Cap.getCapturedVar()->getType().isConstant(CGF.getContext()) &&
           Cap.getCaptureKind() == CapturedStmt::VCK_ByRef)
         return MappableExprsHandler::OMP_MAP_ALWAYS |
                MappableExprsHandler::OMP_MAP_TO;
       if (Cap.getCapturedVar()->getType()->isAnyPointerType())
         return MappableExprsHandler::OMP_MAP_TO |
                MappableExprsHandler::OMP_MAP_PTR_AND_OBJ;
       return MappableExprsHandler::OMP_MAP_PRIVATE |
              MappableExprsHandler::OMP_MAP_TO;
     }
     return MappableExprsHandler::OMP_MAP_TO |
            MappableExprsHandler::OMP_MAP_FROM;
public:
   }
 
   /// Generate all the base pointers, section pointers, sizes and map
   /// types for the extracted mappable expressions. 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(MapBaseValuesArrayTy &BasePointers,
                        MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
-                       MapFlagsArrayTy &Types) const {
+                       MapFlagsArrayTy &Types, MapDimArrayTy &Dims,
+                       MapNonContiguousArrayTy &Offsets,
+                       MapNonContiguousArrayTy &Counts,
+                       MapNonContiguousArrayTy &Strides) 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;
 
     // Helper function to fill the information map for the different supported
     // clauses.
     auto &&InfoGen = [&Info](
for (const auto &M : Info) {
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // Temporary versions of arrays
       MapBaseValuesArrayTy CurBasePointers;
       MapValuesArrayTy CurPointers;
       MapValuesArrayTy CurSizes;
       MapFlagsArrayTy CurTypes;
+      MapDimArrayTy CurDims;
+      MapNonContiguousArrayTy CurOffsets;
+      MapNonContiguousArrayTy CurCounts;
+      MapNonContiguousArrayTy CurStrides;
       StructRangeInfoTy PartialStruct;
 
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
 
         // Remember the current base pointer index.
         unsigned CurrentBasePointersIdx = CurBasePointers.size();
         generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
                                      CurBasePointers, CurPointers, CurSizes,
-                                     CurTypes, PartialStruct,
-                                     IsFirstComponentList, L.IsImplicit);
+                                     CurTypes, CurDims, PartialStruct,
+                                     IsFirstComponentList, L.IsImplicit,
+                                     /*OverlappedElements=*/llvm::None,
+                                     L.Components.back().isNonContiguous(),
+                                     &CurOffsets, &CurCounts, &CurStrides);
 
         // If this entry relates with a device pointer, set the relevant
         // declaration and add the 'return pointer' flag.
         if (L.ReturnDevicePointer) {
           assert(CurBasePointers.size() > CurrentBasePointersIdx &&
                  "Unexpected number of mapped base pointers.");
 
           const ValueDecl *RelevantVD =
               L.Components.back().getAssociatedDeclaration();
           assert(RelevantVD &&
for (const auto &M : Info) {
         emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
                           PartialStruct);
 
       // We need to append the results of this capture to what we already have.
       BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
       Pointers.append(CurPointers.begin(), CurPointers.end());
       Sizes.append(CurSizes.begin(), CurSizes.end());
       Types.append(CurTypes.begin(), CurTypes.end());
+      Dims.append(CurDims.begin(), CurDims.end());
+      Offsets.append(CurOffsets.begin(), CurOffsets.end());
+      Counts.append(CurCounts.begin(), CurCounts.end());
+      Strides.append(CurStrides.begin(), CurStrides.end());
     }
   }
 
   /// Generate all the base pointers, section pointers, sizes and map types for
   /// the extracted map clauses of user-defined mapper.
   void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers,
                                 MapValuesArrayTy &Pointers,
                                 MapValuesArrayTy &Sizes,
for (const auto &M : Info) {
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // Temporary versions of arrays
       MapBaseValuesArrayTy CurBasePointers;
       MapValuesArrayTy CurPointers;
       MapValuesArrayTy CurSizes;
       MapFlagsArrayTy CurTypes;
+      MapDimArrayTy CurDims;
+      MapNonContiguousArrayTy CurOffsets;
+      MapNonContiguousArrayTy CurCounts;
+      MapNonContiguousArrayTy CurStrides;
       StructRangeInfoTy PartialStruct;
 
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
         generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
                                      CurBasePointers, CurPointers, CurSizes,
-                                     CurTypes, PartialStruct,
+                                     CurTypes, CurDims, PartialStruct,
                                      IsFirstComponentList, L.IsImplicit);
         IsFirstComponentList = false;
       }
 
       // 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()) {
+        // Make sure Dims have the same size as BP, P, Sizes, and Types.
+        // Put 0 here to make sure that `emitOffloadingArrays` use it
+        // to skip processing this one. (OpenMP do not allow non-contigous for
+        // declare mapper)
+        CurDims.push_back(0);
         emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
                           PartialStruct);
+      }
 
       // We need to append the results of this capture to what we already have.
       BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
       Pointers.append(CurPointers.begin(), CurPointers.end());
       Sizes.append(CurSizes.begin(), CurSizes.end());
       Types.append(CurTypes.begin(), CurTypes.end());
     }
   }
public:
 
   /// Generate the base pointers, section pointers, sizes and map types
   /// associated to a given capture.
   void generateInfoForCapture(const CapturedStmt::Capture *Cap,
                               llvm::Value *Arg,
                               MapBaseValuesArrayTy &BasePointers,
                               MapValuesArrayTy &Pointers,
                               MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+                              MapDimArrayTy &Dims,
                               StructRangeInfoTy &PartialStruct) const {
     assert(!Cap->capturesVariableArrayType() &&
            "Not expecting to generate map info for a variable array type!");
 
     // We need to know when we generating information for the first component
     const ValueDecl *VD = Cap->capturesThis()
                               ? nullptr
                               : Cap->getCapturedVar()->getCanonicalDecl();
for (const auto &Pair : OverlappedData) {
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
       std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedComponents = Pair.getSecond();
       bool IsFirstComponentList = true;
       generateInfoForComponentList(MapType, MapModifiers, Components,
-                                   BasePointers, Pointers, Sizes, Types,
+                                   BasePointers, Pointers, Sizes, Types, Dims,
                                    PartialStruct, IsFirstComponentList,
                                    IsImplicit, OverlappedComponents);
     }
     // 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;
       std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
       auto It = OverlappedData.find(&L);
       if (It == OverlappedData.end())
-        generateInfoForComponentList(MapType, MapModifiers, Components,
-                                     BasePointers, Pointers, Sizes, Types,
-                                     PartialStruct, IsFirstComponentList,
-                                     IsImplicit);
+        generateInfoForComponentList(
+            MapType, MapModifiers, Components, BasePointers, Pointers, Sizes,
+            Types, Dims, PartialStruct, IsFirstComponentList, IsImplicit);
       IsFirstComponentList = false;
     }
   }
 
   /// Generate the base pointers, section pointers, sizes and map types
   /// associated with the declare target link variables.
   void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers,
                                         MapValuesArrayTy &Pointers,
                                         MapValuesArrayTy &Sizes,
-                                        MapFlagsArrayTy &Types) const {
+                                        MapFlagsArrayTy &Types,
+                                        MapDimArrayTy &Dims) const {
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
     // Map other list items in the map clause which are not captured variables
     // but "declare target link" global variables.
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
       for (const auto L : C->component_lists()) {
         if (!L.first)
           continue;
         const auto *VD = dyn_cast<VarDecl>(L.first);
         if (!VD)
           continue;
         llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
             OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
         if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
             !Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
           continue;
         StructRangeInfoTy PartialStruct;
         generateInfoForComponentList(
             C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
-            Pointers, Sizes, Types, PartialStruct,
+            Pointers, Sizes, Types, Dims, PartialStruct,
             /*IsFirstComponentList=*/true, C->isImplicit());
         assert(!PartialStruct.Base.isValid() &&
                "No partial structs for declare target link expected.");
       }
     }
   }
 
   /// Generate the default map information for a given capture \a CI,
void generateDefaultMapInfo(const CapturedStmt::Capture &CI,
 
     // Add flag stating this is an implicit map.
     if (IsImplicit)
       CurMapTypes.back() |= OMP_MAP_IMPLICIT;
   }
 };
 } // anonymous namespace
 
 /// Emit the arrays used to pass the captures and map information to the
 /// offloading runtime library. If there is no map or capture information,
 /// return nullptr by reference.
 static void
 emitOffloadingArrays(CodeGenFunction &CGF,
                      MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
                      MappableExprsHandler::MapValuesArrayTy &Pointers,
                      MappableExprsHandler::MapValuesArrayTy &Sizes,
                      MappableExprsHandler::MapFlagsArrayTy &MapTypes,
-                     CGOpenMPRuntime::TargetDataInfo &Info) {
+                     MappableExprsHandler::MapDimArrayTy &Dims,
+                     CGOpenMPRuntime::TargetDataInfo &Info,
+                     bool IsNonContiguous = false) {
   CodeGenModule &CGM = CGF.CGM;
   ASTContext &Ctx = CGF.getContext();
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = BasePointers.size();
 
   if (Info.NumberOfPtrs) {
     // Detect if we have any capture size requiring runtime evaluation of the
     // size so that a constant array could be eventually used.
     bool hasRuntimeEvaluationCaptureSize = false;
     for (llvm::Value *S : Sizes)
       if (!isa<llvm::Constant>(S)) {
         hasRuntimeEvaluationCaptureSize = true;
         break;
       }
if (hasRuntimeEvaluationCaptureSize) {
           Int64Ty, PointerNumAP, nullptr, ArrayType::Normal,
           /*IndexTypeQuals=*/0);
       Info.SizesArray =
           CGF.CreateMemTemp(SizeArrayType, ".offload_sizes").getPointer();
     } else {
       // We expect all the sizes to be constant, so we collect them to create
       // a constant array.
       SmallVector<llvm::Constant *, 16> ConstSizes;
-      for (llvm::Value *S : Sizes)
-        ConstSizes.push_back(cast<llvm::Constant>(S));
+      for (unsigned I = 0, E = Sizes.size(); I < E; ++I) {
+        if (IsNonContiguous &&
+            (MapTypes[I] & MappableExprsHandler::OMP_MAP_DESCRIPTOR)) {
+          ConstSizes.push_back(llvm::ConstantInt::get(CGF.Int64Ty, Dims[I]));
+        } else {
+          ConstSizes.push_back(cast<llvm::Constant>(Sizes[I]));
+        }
+      }
 
       auto *SizesArrayInit = llvm::ConstantArray::get(
           llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes);
       std::string Name = CGM.getOpenMPRuntime().getName({"offload_sizes"});
       auto *SizesArrayGbl = new llvm::GlobalVariable(
           CGM.getModule(), SizesArrayInit->getType(),
           /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
           SizesArrayInit, Name);
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
             /*Idx1=*/I);
         Address SAddr(S, Ctx.getTypeAlignInChars(Int64Ty));
         CGF.Builder.CreateStore(
             CGF.Builder.CreateIntCast(Sizes[I], CGM.Int64Ty, /*isSigned=*/true),
             SAddr);
       }
     }
   }
+
+  if (IsNonContiguous) {
+    if (Info.Offsets.empty())
+      return;
+
+    ASTContext &C = CGF.getContext();
+    CodeGenModule &CGM = CGF.CGM;
+
+    // Build an array of struct descriptor_dim and then assign it to
+    // offload_args.
+    if (Info.NumberOfPtrs) {
+      // Build struct descriptor_dim {
+      //  int64_t offset;
+      //  int64_t count;
+      //  int64_t stride
+      // };
+      QualType Int64Ty =
+          C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/true);
+      RecordDecl *RD;
+      RD = C.buildImplicitRecord("descriptor_dim");
+      RD->startDefinition();
+      addFieldToRecordDecl(C, RD, Int64Ty);
+      addFieldToRecordDecl(C, RD, Int64Ty);
+      addFieldToRecordDecl(C, RD, Int64Ty);
+      RD->completeDefinition();
+      QualType DimTy = C.getRecordType(RD);
+
+      enum { OffsetFD = 0, CountFD, StrideFD };
+      // The reason we need two index variable here is because the size of
+      // "Dims" is the same as the size of Components, however, the size of
+      // offset, count , and stride is equal to the size of base declaration
+      // that is non-contiguous.
+      for (unsigned I = 0, L = 0, E = Info.Offsets.size(); I < E; ++I) {
+        if (Dims[I] == 0)
+          continue;
+        llvm::APInt Size(/*numBits=*/32, Dims[I]);
+        QualType ArrayTy =
+            C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0);
+        Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims");
+        for (unsigned II = 0, EE = Dims[I]; II < EE; ++II) {
+          unsigned RevIdx = EE - II - 1;
+          LValue DimsLVal = CGF.MakeAddrLValue(
+              CGF.Builder.CreateConstArrayGEP(DimsAddr, II), DimTy);
+          // Offset
+          LValue OffsetLVal = CGF.EmitLValueForField(
+              DimsLVal, *std::next(RD->field_begin(), OffsetFD));
+          CGF.EmitStoreOfScalar(Info.Offsets[L][RevIdx], OffsetLVal);
+          // Count
+          LValue CountLVal = CGF.EmitLValueForField(
+              DimsLVal, *std::next(RD->field_begin(), CountFD));
+          CGF.EmitStoreOfScalar(Info.Counts[L][RevIdx], CountLVal);
+          // Stride
+          LValue StrideLVal = CGF.EmitLValueForField(
+              DimsLVal, *std::next(RD->field_begin(), StrideFD));
+          CGF.EmitStoreOfScalar(Info.Strides[L][RevIdx], StrideLVal);
+        }
+        // args[I] = &dims
+        Address DAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+            DimsAddr, CGM.Int8PtrTy);
+        llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(
+            llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+            Info.PointersArray, 0, I);
+        Address PAddr(P, C.getTypeAlignInChars(C.VoidPtrTy));
+        CGF.Builder.CreateStore(DAddr.getPointer(), PAddr);
+        ++L;
+      }
+    }
+  }
 }
 
 /// Emit the arguments to be passed to the runtime library based on the
 /// arrays of pointers, sizes and map types.
 static void emitOffloadingArraysArgument(
     CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
     llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
     llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) {
   CodeGenModule &CGM = CGF.CGM;
   if (Info.NumberOfPtrs) {
     BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.BasePointersArray,
         /*Idx0=*/0, /*Idx1=*/0);
void CGOpenMPRuntime::emitTargetCall(
   auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
                           &CapturedVars, RequiresOuterTask,
                           &CS](CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the captured variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapDimArrayTy Dims;
 
     // Get mappable expression information.
     MappableExprsHandler MEHandler(D, CGF);
     llvm::DenseMap<llvm::Value *, llvm::Value *> LambdaPointers;
 
     auto RI = CS.getCapturedRecordDecl()->field_begin();
     auto CV = CapturedVars.begin();
     for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
                                               CE = CS.capture_end();
          CI != CE; ++CI, ++RI, ++CV) {
       MappableExprsHandler::MapBaseValuesArrayTy CurBasePointers;
       MappableExprsHandler::MapValuesArrayTy CurPointers;
       MappableExprsHandler::MapValuesArrayTy CurSizes;
       MappableExprsHandler::MapFlagsArrayTy CurMapTypes;
+      MappableExprsHandler::MapDimArrayTy CurDims;
       MappableExprsHandler::StructRangeInfoTy PartialStruct;
 
       // VLA sizes are passed to the outlined region by copy and do not have map
       // information associated.
       if (CI->capturesVariableArrayType()) {
         CurBasePointers.push_back(*CV);
         CurPointers.push_back(*CV);
         CurSizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(RI->getType()), CGF.Int64Ty, /*isSigned=*/true));
         // Copy to the device as an argument. No need to retrieve it.
         CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_LITERAL |
                               MappableExprsHandler::OMP_MAP_TARGET_PARAM |
                               MappableExprsHandler::OMP_MAP_IMPLICIT);
       } else {
         // If we have any information in the map clause, we use it, otherwise we
         // just do a default mapping.
         MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
-                                         CurSizes, CurMapTypes, PartialStruct);
+                                         CurSizes, CurMapTypes, CurDims,
+                                         PartialStruct);
         if (CurBasePointers.empty())
           MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
                                            CurPointers, CurSizes, CurMapTypes);
         // Generate correct mapping for variables captured by reference in
         // lambdas.
         if (CI->capturesVariable())
           MEHandler.generateInfoForLambdaCaptures(
               CI->getCapturedVar(), *CV, CurBasePointers, CurPointers, CurSizes,
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
       MapTypes.append(CurMapTypes.begin(), CurMapTypes.end());
     }
     // Adjust MEMBER_OF flags for the lambdas captures.
     MEHandler.adjustMemberOfForLambdaCaptures(LambdaPointers, BasePointers,
                                               Pointers, MapTypes);
     // Map other list items in the map clause which are not captured variables
     // but "declare target link" global variables.
     MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes,
-                                               MapTypes);
+                                               MapTypes, Dims);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims,
+                         Info);
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
                                  Info.MapTypesArray, Info);
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());
     InputInfo.PointersArray =
         Address(Info.PointersArray, CGM.getPointerAlign());
void CGOpenMPRuntime::emitTargetDataCalls(
   // closing of the region.
   auto &&BeginThenGen = [this, &D, Device, &Info,
                          &CodeGen](CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the mapped variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapDimArrayTy Dims;
+    MappableExprsHandler::MapNonContiguousArrayTy Offsets;
+    MappableExprsHandler::MapNonContiguousArrayTy Counts;
+    MappableExprsHandler::MapNonContiguousArrayTy Strides;
 
     // Get map clause information.
     MappableExprsHandler MCHandler(D, CGF);
-    MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+    MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Dims,
+                              Offsets, Counts, Strides);
+
+    // Fill up non-contiguous information.
+    Info.Offsets = std::move((Offsets));
+    Info.Counts = std::move((Counts));
+    Info.Strides = std::move((Strides));
 
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims,
+                         Info, /*IsNonContiguous=*/true);
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
                                  SizesArrayArg, MapTypesArrayArg, Info);
 
void CGOpenMPRuntime::emitTargetDataStandAloneCall(
 
   auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](
                              CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the mapped variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapDimArrayTy Dims;
+    MappableExprsHandler::MapNonContiguousArrayTy Offsets;
+    MappableExprsHandler::MapNonContiguousArrayTy Counts;
+    MappableExprsHandler::MapNonContiguousArrayTy Strides;
 
     // Get map clause information.
     MappableExprsHandler MEHandler(D, CGF);
-    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Dims,
+                              Offsets, Counts, Strides);
 
     TargetDataInfo Info;
+
+    // Fill up non-contiguous information.
+    Info.Offsets = std::move((Offsets));
+    Info.Counts = std::move((Counts));
+    Info.Strides = std::move((Strides));
+
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims,
+                         Info, /*IsNonContiguous=*/true);
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
                                  Info.MapTypesArray, Info);
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());
     InputInfo.PointersArray =
         Address(Info.PointersArray, CGM.getPointerAlign());

clang/lib/Sema/SemaOpenMP.cpp

 
 //===----------------------------------------------------------------------===//
 // Stack of data-sharing attributes for variables
 //===----------------------------------------------------------------------===//
 
 static const Expr *checkMapClauseExpressionBase(
     Sema &SemaRef, Expr *E,
     OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
-    OpenMPClauseKind CKind, bool NoDiagnose);
+    OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose);
 
 namespace {
 /// Default data sharing attributes, which can be applied to directive.
 enum DefaultDataSharingAttributes {
   DSA_unspecified = 0, /// Data sharing attribute not specified.
   DSA_none = 1 << 0,   /// Default data sharing attribute 'none'.
   DSA_shared = 1 << 1, /// Default data sharing attribute 'shared'.
 };
if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParenCasts())) {
         if (DVar.CKind != OMPC_unknown)
           ImplicitFirstprivate.push_back(E);
       }
       return;
     }
     if (isOpenMPTargetExecutionDirective(DKind)) {
       OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
       if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+                                        Stack->getCurrentDirective(),
                                         /*NoDiagnose=*/true))
         return;
       const auto *VD = cast<ValueDecl>(
           CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl());
       if (!Stack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&CurComponents](
                   OMPClauseMappableExprCommon::MappableExprComponentListRef

 // but these would be valid:
 //   r.ArrS[3].Arr[6:7]
 //
 //   r.ArrS[3].x
 namespace {
 class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
   Sema &SemaRef;
   OpenMPClauseKind CKind = OMPC_unknown;
+  OpenMPDirectiveKind DKind = OMPD_unknown;
   OMPClauseMappableExprCommon::MappableExprComponentList &Components;
+  bool IsNonContiguous = false;
   bool NoDiagnose = false;
   const Expr *RelevantExpr = nullptr;
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
+  bool AllowAnotherPtr = true;
   SourceLocation ELoc;
   SourceRange ERange;
 
   void emitErrorMsg() {
     // If nothing else worked, this is not a valid map clause expression.
     if (SemaRef.getLangOpts().OpenMP < 50) {
       SemaRef.Diag(ELoc,
                    diag::err_omp_expected_named_var_member_or_array_expression)
           << ERange;
     } else {
       SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses)
           << getOpenMPClauseName(CKind) << ERange;
     }
   }
 
 public:
   bool VisitDeclRefExpr(DeclRefExpr *DRE) {
     if (!isa<VarDecl>(DRE->getDecl())) {
       emitErrorMsg();
       return false;
     }
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     RelevantExpr = DRE;
     // Record the component.
-    Components.emplace_back(DRE, DRE->getDecl());
+    Components.emplace_back(DRE, DRE->getDecl(), IsNonContiguous);
     return true;
   }
 
   bool VisitMemberExpr(MemberExpr *ME) {
     Expr *E = ME;
     Expr *BaseE = ME->getBase()->IgnoreParenCasts();
 
     if (isa<CXXThisExpr>(BaseE)) {
bool VisitMemberExpr(MemberExpr *ME) {
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
     //  If a list item is an element of a structure, only the rightmost symbol
     //  of the variable reference can be an array section.
     //
     AllowUnitySizeArraySection = false;
     AllowWholeSizeArraySection = false;
 
     // Record the component.
-    Components.emplace_back(ME, FD);
+    Components.emplace_back(ME, FD, IsNonContiguous);
     return RelevantExpr || Visit(E);
   }
 
   bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
     Expr *E = AE->getBase()->IgnoreParenImpCasts();
 
     if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
       if (!NoDiagnose) {
if (const auto *TE = dyn_cast<CXXThisExpr>(E->IgnoreParenCasts())) {
         SemaRef.Diag(AE->getIdx()->getExprLoc(),
                      diag::note_omp_invalid_subscript_on_this_ptr_map);
       }
       assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(AE, nullptr);
+    Components.emplace_back(AE, nullptr, IsNonContiguous);
 
     return RelevantExpr || Visit(E);
   }
 
   bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
     assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
     Expr *E = OASE->getBase()->IgnoreParenImpCasts();
     QualType CurType =
if (AllowWholeSizeArraySection) {
       // Any array section is currently allowed. Allowing a whole size array
       // section implies allowing a unity array section as well.
       //
       // If this array section refers to the whole dimension we can still
       // accept other array sections before this one, except if the base is a
       // pointer. Otherwise, only unitary sections are accepted.
       if (NotWhole || IsPointer)
         AllowWholeSizeArraySection = false;
+    } else if (DKind == OMPD_target_update &&
+               SemaRef.getLangOpts().OpenMP >= 50) {
+      if (IsPointer && !AllowAnotherPtr)
+        SemaRef.Diag(ELoc, diag::err_omp_section_length_undefined) << true;
+      else
+        IsNonContiguous = true;
     } else if (AllowUnitySizeArraySection && NotUnity) {
       // A unity or whole array section is not allowed and that is not
       // compatible with the properties of the current array section.
       SemaRef.Diag(
         ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
         << OASE->getSourceRange();
       return false;
     }
 
+    if (IsPointer)
+      AllowAnotherPtr = false;
+
     if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
       Expr::EvalResult ResultR;
       Expr::EvalResult ResultL;
       if (!OASE->getLength()->isValueDependent() &&
           OASE->getLength()->EvaluateAsInt(ResultR, SemaRef.getASTContext()) &&
           !ResultR.Val.getInt().isOneValue()) {
         SemaRef.Diag(OASE->getLength()->getExprLoc(),
                      diag::err_omp_invalid_map_this_expr);
if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
         SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
                      diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
       }
       assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(OASE, nullptr);
+    Components.emplace_back(OASE, nullptr, false);
     return RelevantExpr || Visit(E);
   }
   bool VisitOMPArrayShapingExpr(OMPArrayShapingExpr *E) {
     Expr *Base = E->getBase();
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(E, nullptr);
+    Components.emplace_back(E, nullptr, IsNonContiguous);
 
     return Visit(Base->IgnoreParenImpCasts());
   }
 
   bool VisitUnaryOperator(UnaryOperator *UO) {
     if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() ||
         UO->getOpcode() != UO_Deref) {
       emitErrorMsg();
       return false;
     }
     if (!RelevantExpr) {
       // Record the component if haven't found base decl.
-      Components.emplace_back(UO, nullptr);
+      Components.emplace_back(UO, nullptr, false);
     }
     return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
   }
   bool VisitBinaryOperator(BinaryOperator *BO) {
     if (SemaRef.getLangOpts().OpenMP < 50 || !BO->getType()->isPointerType()) {
       emitErrorMsg();
       return false;
     }
 
     // Pointer arithmetic is the only thing we expect to happen here so after we
     // make sure the binary operator is a pointer type, the we only thing need
     // to to is to visit the subtree that has the same type as root (so that we
     // know the other subtree is just an offset)
     Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
     Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
-    Components.emplace_back(BO, nullptr);
+    Components.emplace_back(BO, nullptr, false);
     assert((LE->getType().getTypePtr() == BO->getType().getTypePtr() ||
             RE->getType().getTypePtr() == BO->getType().getTypePtr()) &&
            "Either LHS or RHS have base decl inside");
     if (BO->getType().getTypePtr() == LE->getType().getTypePtr())
       return RelevantExpr || Visit(LE);
     return RelevantExpr || Visit(RE);
   }
   bool VisitCXXThisExpr(CXXThisExpr *CTE) {
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     RelevantExpr = CTE;
-    Components.emplace_back(CTE, nullptr);
+    Components.emplace_back(CTE, nullptr, IsNonContiguous);
     return true;
   }
   bool VisitStmt(Stmt *) {
     emitErrorMsg();
     return false;
   }
   const Expr *getFoundBase() const {
     return RelevantExpr;
   }
   explicit MapBaseChecker(
-      Sema &SemaRef, OpenMPClauseKind CKind,
+      Sema &SemaRef, OpenMPClauseKind CKind, OpenMPDirectiveKind DKind,
       OMPClauseMappableExprCommon::MappableExprComponentList &Components,
       bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange)
-      : SemaRef(SemaRef), CKind(CKind), Components(Components),
+      : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components),
         NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
 };
 } // namespace
 
 /// Return the expression of the base of the mappable expression or null if it
 /// cannot be determined and do all the necessary checks to see if the expression
 /// is valid as a standalone mappable expression. In the process, record all the
 /// components of the expression.
 static const Expr *checkMapClauseExpressionBase(
     Sema &SemaRef, Expr *E,
     OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
-    OpenMPClauseKind CKind, bool NoDiagnose) {
+    OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
-  MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
+  MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents, NoDiagnose, ELoc,
                          ERange);
   if (Checker.Visit(E->IgnoreParens()))
     return Checker.getFoundBase();
   return nullptr;
 }
 
 // Return true if expression E associated with value VD has conflicts with other
 // map information.
for (Expr *RE : MVLI.VarList) {
     }
 
     OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
     ValueDecl *CurDeclaration = nullptr;
 
     // Obtain the array or member expression bases if required. Also, fill the
     // components array with all the components identified in the process.
     const Expr *BE = checkMapClauseExpressionBase(
-        SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false);
+        SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(),
+        /*NoDiagnose=*/false);
     if (!BE)
       continue;
 
     assert(!CurComponents.empty() &&
            "Invalid mappable expression information.");
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) {
       // Add store "this" pointer to class in DSAStackTy for future checking
for (Expr *RefExpr : VarList) {
     // We need to add a data sharing attribute for this variable to make sure it
     // is correctly captured. A variable that shows up in a use_device_ptr has
     // similar properties of a first private variable.
     DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
 
     // Create a mappable component for the list item. List items in this clause
     // only need a component.
     MVLI.VarBaseDeclarations.push_back(D);
     MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
-    MVLI.VarComponents.back().push_back(
-        OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+    MVLI.VarComponents.back().emplace_back(
+        OMPClauseMappableExprCommon::MappableComponent(
+            SimpleRefExpr, D,
+            /*IsNonContiguous=*/false));
   }
 
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDevicePtrClause::Create(
       Context, Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
       MVLI.VarBaseDeclarations, MVLI.VarComponents);
 }
for (Expr *RefExpr : VarList) {
     // is correctly captured. A variable that shows up in a use_device_addr has
     // similar properties of a first private variable.
     DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
 
     // Create a mappable component for the list item. List items in this clause
     // only need a component.
     MVLI.VarBaseDeclarations.push_back(D);
     MVLI.VarComponents.emplace_back();
-    MVLI.VarComponents.back().push_back(
-        OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+    MVLI.VarComponents.back().emplace_back(
+        OMPClauseMappableExprCommon::MappableComponent(
+            SimpleRefExpr, D,
+            /*IsNonContiguous=*/false));
   }
 
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDeviceAddrClause::Create(Context, Locs, MVLI.ProcessedVarList,
                                         MVLI.VarBaseDeclarations,
                                         MVLI.VarComponents);
 }
if (DSAStack->checkMappableExprComponentListsForDecl(
       Diag(ELoc, diag::err_omp_map_shared_storage) << RefExpr->getSourceRange();
       Diag(ConflictExpr->getExprLoc(), diag::note_used_here)
           << ConflictExpr->getSourceRange();
       continue;
     }
 
     // Store the components in the stack so that they can be used to check
     // against other clauses later on.
-    OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D);
+    OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D, false);
     DSAStack->addMappableExpressionComponents(
         D, MC, /*WhereFoundClauseKind=*/OMPC_is_device_ptr);
 
     // Record the expression we've just processed.
     MVLI.ProcessedVarList.push_back(SimpleRefExpr);
 
     // Create a mappable component for the list item. List items in this clause
     // only need a component. We use a null declaration to signal fields in

clang/lib/Serialization/ASTReader.cpp

void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readExpr();
+    Expr *AssociatedExprPr = Record.readExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExprPr, AssociatedDecl, /*IsNonContiguous=*/false));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPAllocateClause(OMPAllocateClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   C->setColonLoc(Record.readSourceLocation());
   C->setAllocator(Record.readSubExpr());
void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    Expr *AssociatedExprPr = Record.readSubExpr();
+    bool IsNonContiguous = Record.readBool();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExprPr, AssociatedDecl, IsNonContiguous));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
   C->setMapperIdInfo(Record.readDeclarationNameInfo());
   auto NumVars = C->varlist_size();
void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    Expr *AssociatedExprPr = Record.readSubExpr();
+    bool IsNonContiguous = Record.readBool();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExprPr, AssociatedDecl, IsNonContiguous));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    auto *AssociatedExprPr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExprPr, AssociatedDecl, /*IsNonContiguous=*/false));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
     Expr *AssociatedExpr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExpr, AssociatedDecl, /*IsNonContiguous*/ false));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
   auto TotalLists = C->getTotalComponentListNum();
for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
     Expr *AssociatedExpr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExpr, AssociatedDecl, /*IsNonContiguous=*/false));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPNontemporalClause(OMPNontemporalClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   unsigned NumVars = C->varlist_size();
   SmallVector<Expr *, 16> Vars;
   Vars.reserve(NumVars);

clang/lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())
     Record.push_back(N);
   for (auto N : C->all_lists_sizes())
     Record.push_back(N);
   for (auto &M : C->all_components()) {
     Record.AddStmt(M.getAssociatedExpression());
+    Record.writeBool(M.isNonContiguous());
     Record.AddDeclRef(M.getAssociatedDeclaration());
   }
 }
 
 void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());
   Record.push_back(C->getTotalComponentsNum());
   Record.AddSourceLocation(C->getLParenLoc());
   Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
   Record.AddDeclarationNameInfo(C->getMapperIdInfo());
   for (auto *E : C->varlists())
     Record.AddStmt(E);
   for (auto *E : C->mapperlists())
     Record.AddStmt(E);
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())
     Record.push_back(N);
   for (auto N : C->all_lists_sizes())
     Record.push_back(N);
   for (auto &M : C->all_components()) {
     Record.AddStmt(M.getAssociatedExpression());
+    Record.writeBool(M.isNonContiguous());
     Record.AddDeclRef(M.getAssociatedDeclaration());
   }
 }
 
 void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());

clang/test/OpenMP/target_update_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void foo() {}
 
 template <class T, class U>
 T foo(T targ, U uarg) {
   static T a, *p;
   U b;
   int l;
 #pragma omp target update to(([a][targ])p, a) if(l>5) device(l) nowait depend(inout:l)
 
 #pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l)
+
+#ifdef OMP5
+  U marr[10][10][10];
+#pragma omp target update to(marr[2] [0:2] [0:2])
+
+#pragma omp target update from(marr[2] [0:2] [0:2])
+
+#pragma omp target update to(marr[:] [0:2] [0:2])
+
+#pragma omp target update from(marr[:] [0:2] [0:2])
+
+#pragma omp target update to(marr[:][:l] [l:])
+
+#pragma omp target update from(marr[:][:l] [l:])
+
+#pragma omp target update to(marr[:2][:1][:])
+
+#pragma omp target update from(marr[:2][:1][:])
+
+#pragma omp target update to(marr[:2][:][:1])
+
+#pragma omp target update from(marr[:2][:][:1])
+
+#pragma omp target update to(marr[:2][:] [1:])
+
+#pragma omp target update from(marr[:2][:] [1:])
+
+#pragma omp target update to(marr[:1] [3:2][:2])
+
+#pragma omp target update from(marr[:1] [3:2][:2])
+
+#pragma omp target update to(marr[:1][:2][0])
+
+#pragma omp target update from(marr[:1][:2][0])
+
+// OMP5: marr[10][10][10];
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][:l][l:])
+// OMP5-NEXT: #pragma omp target update from(marr[:][:l][l:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+#endif
+
   return a + targ + (T)b;
 }
 // CHECK:      static T a, *p;
 // CHECK-NEXT: U b;
-// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l){{$}}
-// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
-// CHECK:      static int a, *p;
-// CHECK-NEXT: float b;
-// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
-// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
-// CHECK:      static char a, *p;
-// CHECK-NEXT: float b;
-// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
-// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
 
 int main(int argc, char **argv) {
   static int a;
   int n;
   float f;
 
 // CHECK:      static int a;
 // CHECK-NEXT: int n;
 // CHECK-NEXT: float f;
 #pragma omp target update to(a) if(f>0.0) device(n) nowait depend(in:n)
 // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n)
 #pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n)
 // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n)
+
+#ifdef OMP5
+float marr[10][10][10];
+// OMP5: marr[10][10][10];
+#pragma omp target update to(marr[2] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+#pragma omp target update from(marr[2] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+#pragma omp target update to(marr[:] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+#pragma omp target update from(marr[:] [0:2] [0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+#pragma omp target update to(marr[:][:n] [n:])
+// OMP5: #pragma omp target update to(marr[:][:n][n:])
+#pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+#pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+#pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+#pragma omp target update to(marr[:2][:] [1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+#pragma omp target update from(marr[:2][:] [1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+#pragma omp target update to(marr[:1] [3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+#pragma omp target update from(marr[:1] [3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+#pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+#pragma omp target update from(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+#endif
+
   return foo(argc, f) + foo(argv[0][0], f) + a;
 }
 
 #endif

clang/test/OpenMP/target_update_codegen.cpp

void array_shaping(float *f, int sa) {
   // CK18-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
   // CK18-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
   // CK18-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
   // CK18-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
   #pragma omp target update from(([sa][5])f)
 }
 
 #endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=50 -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=50 -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 CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK19 --check-prefix CK19-32
+// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=50 -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=50 -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 CK19 --check-prefix CK19-32
+
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK19
+
+// CK19: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK19: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3]
+// CK19: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK19-LABEL: _Z3foo
+void foo(int arg) {
+  int arr[3][4][5];
+
+  // CK19: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+  // CK19: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x i32]]], [3 x [4 x [5 x i32]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK19: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x i32]], [4 x [5 x i32]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK19: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_DECAY]], {{.+}}
+  // CK19: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+  // CK19: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 1
+  // CK19: [[LEN:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]]
+  // CK19: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK19: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK19: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK19: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK19: store i64 0, i64* [[OFFSET]],
+  // CK19: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK19: store i64 2, i64* [[COUNT]],
+  // CK19: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK19: store i64 80, i64* [[STRIDE]],
+  // CK19: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK19: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK19: store i64 [[ARG:%.+]], i64* [[OFFSET_2]],
+  // CK19: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK19: store i64 [[LEN]], i64* [[COUNT_2]],
+  // CK19: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK19: store i64 20, i64* [[STRIDE_2]],
+  // CK19: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+  // CK19: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+  // CK19: store i64 1, i64* [[OFFSET_3]],
+  // CK19: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+  // CK19: store i64 4, i64* [[COUNT_3]],
+  // CK19: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+  // CK19: store i64 4, i64* [[STRIDE_3]],
+
+  // CK19-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK19-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK19-DAG: store i8* [[PC0]], i8** [[PTRS]],
+
+#pragma omp target update to(arr [0:2] [arg:] [1:4])
+  { ++arg; }
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -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=50 -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 CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK20 --check-prefix CK20-32
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -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=50 -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 CK20 --check-prefix CK20-32
+
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK20
+
+struct ST {
+  int a;
+  double *b;
+};
+
+// CK20: [[STRUCT_ST:%.+]] = type { i32, double* }
+// CK20: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK20: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 2]
+// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK20-LABEL: _Z3foo
+void foo(int arg) {
+  ST arr[3][4];
+  // CK20: [[DIMS:%.+]] = alloca [2 x [[STRUCT_DESCRIPTOR]]],
+  // CK20: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [[STRUCT_ST]]]], [3 x [4 x [[STRUCT_ST]]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [[STRUCT_ST]]], [4 x [[STRUCT_ST]]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK20: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [[STRUCT_ST]], [[STRUCT_ST]]* [[ARRAY_DECAY]], {{.+}}
+  // CK20: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [[STRUCT_ST]]]]**
+  // CK20: store [3 x [4 x [[STRUCT_ST]]]]* [[ARR]], [3 x [4 x [[STRUCT_ST]]]]** [[BPC]],
+  // CK20: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[PC:%.+]] = bitcast i8** [[P0]] to [[STRUCT_ST]]**
+  // CK20: store [[STRUCT_ST]]* [[ARRAY_IDX_1]], [[STRUCT_ST]]** [[PC]],
+  // CK20: [[DIM_1:%.+]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR]]], [2 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK20: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK20: store i64 0, i64* [[OFFSET]],
+  // CK20: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK20: store i64 2, i64* [[COUNT]],
+  // CK20: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK20: store i64 {{32|64}}, i64* [[STRIDE]],
+  // CK20: [[DIM_2:%.+]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR]]], [2 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK20: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK20: store i64 1, i64* [[OFFSET_2]],
+  // CK20: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK20: store i64 4, i64* [[COUNT_2]],
+  // CK20: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK20: store i64 {{8|16}}, i64* [[STRIDE_2]],
+  // CK20-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK20-DAG: [[PC0:%.+]] = bitcast [2 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK20-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK20-DAG: store i8* [[PC0]], i8** [[PTRS]],
+
+#pragma omp target update to(arr [0:2] [1:4])
+  { ++arg; }
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -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=50 -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 CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK21 --check-prefix CK21-32
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -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=50 -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 CK21 --check-prefix CK21-32
+
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK21
+
+// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x double*]]] }
+// CK21: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 299067162755073]
+
+struct ST {
+  double *dptr[10][10][10];
+
+  // CK21: _ZN2ST3fooEv
+  void foo() {
+    // CK21: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+    // CK21: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x [10 x [10 x double*]]], [10 x [10 x [10 x double*]]]* [[DPTR:%.+]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x [10 x double*]], [10 x [10 x double*]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_DECAY]], {{.+}} 1
+    // CK21: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 0
+    // CK21: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+    // CK21: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK21: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+    // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 0, i64* [[OFFSET]],
+    // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 2, i64* [[COUNT]],
+    // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{400|800}}, i64* [[STRIDE]],
+    // CK21: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+    // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 1, i64* [[OFFSET]],
+    // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 3, i64* [[COUNT]],
+    // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{40|80}}, i64* [[STRIDE]],
+    // CK21: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+    // CK21: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 0, i64* [[OFFSET]],
+    // CK21: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 4, i64* [[COUNT]],
+    // CK21: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{4|8}}, i64* [[STRIDE]],
+    // CK21-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}}* [[GEPSZ:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+    // CK21-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+    // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i32 0, i32 0
+    // CK21-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(dptr [0:2] [1:3] [0:4])
+  }
+};
+
+void bar() {
+  ST st;
+  st.foo();
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -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=50 -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 CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK22 --check-prefix CK22-32
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -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=50 -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 CK22 --check-prefix CK22-32
+
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -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=50 -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-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK22
+
+// CK22: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK22: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3]
+// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+struct ST {
+  // CK22: _ZN2ST3fooEPA10_Pi
+  void foo(int *arr[5][10]) {
+    // CK22: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+    // CK22: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARR:%.+]], {{.+}} 0
+    // CK22: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+    // CK22: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds i32*, i32** [[ARRAY_DECAY:%.+]], {{.+}} 1
+    // CK22: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+    // CK22: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK22: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 0, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 2, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 {{200|400}}, i64* [[STRIDE]],
+    // CK22: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 1, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 3, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 {{40|80}}, i64* [[STRIDE]],
+    // CK22: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 0, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 4, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 {{4|8}}, i64* [[STRIDE]],
+    // CK22-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+    // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+    // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+    // CK22-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+    // CK22-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+    // CK22-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(arr [0:2] [1:3] [0:4])
+  }
+};
+
+void bar() {
+  ST st;
+  int *arr[5][10];
+  st.foo(arr);
+}
+
+#endif
 #endif

clang/test/OpenMP/target_update_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized
 
 void xxx(int argc) {
   int x; // expected-note {{initialize the variable 'x' to silence this warning}}
 #pragma omp target update to(x)
   argc = x; // expected-warning {{variable 'x' is uninitialized when used here}}
 }
 
 void foo() {
int main(int argc, char **argv) {
   #pragma omp target update to(m) [ // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
   #pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
   #pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
 
   #pragma omp target update from(m) allocate(m) // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp target update'}}
   {
     foo();
   }
+
+  double marr[10][5][10];
+#pragma omp target update to(marr [0:] [2:4] [1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(marr [0:] [2:4] [1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  int arr[4][3][2][1];
+#pragma omp target update to(arr [0:2] [2:4][:2][1]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(arr [0:2] [2:4][:2][1]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  double ***dptr;
+#pragma omp target update to(dptr [0:2] [2:4] [1:2]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(dptr [0:2] [2:4] [1:2]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
   return tmain(argc, argv);
 }

clang/test/OpenMP/target_update_to_messages.cpp

#pragma omp target update to(*this) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
     {}
 #pragma omp target update to(*(this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this->S->i+this->S->p)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this->S->i+this->S->s6[0].pp)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(*(this->ptr)+a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this+this)) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
     {}
+
+    double marr[10][5][10];
+#pragma omp target update to(marr [0:] [2:4] [1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
   }
 };
 
 S3 h;
 #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
 
 typedef int from;
 

 #pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update to(s7.p[:10])
 #pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
 #pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update to(s7.x)
   }
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;
   const int da[5] = { 0 };
   S4 e(4);
   S5 g(5);
   int i, t[20];