[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

5315

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

5315

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

16537

@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

8397–8400

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

16537

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

8397–8400

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

clang/lib/Sema/SemaOpenMP.cpp

16537

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

16537

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

4674–4675

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

4686

isNonContiguous()

clang/lib/Sema/SemaOpenMP.cpp

16359

Add default initializer

16539–16543

Remove braces here, they are not needed.

clang/lib/Serialization/ASTWriter.cpp

6575

There is a member function writeBool

6601

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

4679–4680

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

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7891–7892

Restore original formatting

8377–8379

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

8398

Use prefix form ++DimSize.

8457

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

8512

Same, comment not required

8567–8581

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.

9353–9355

Why removed the comment?

9694

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

clang/lib/Sema/SemaOpenMP.cpp

18388–18390

Use .emplace_back(SimpleRefExpr, D, false);

clang/lib/Serialization/ASTReader.cpp

12481–12482

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

12598–12599

Same, use emplace_back()

12655

Same, use emplace_back()

12712

Same, use emplace_back()

12756

Same, use emplace_back()

clang/test/OpenMP/target_update_to_messages.cpp

147 ↗

(On Diff #268009)

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

7823–7826

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?).

8579–8580

Use range-based loop, if possible.

8582–8584

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

8641

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

8646

Same question about array subscript expressions.

9030–9042

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);`

9580–9583

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

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8107–8110

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?

8578

Expand auto here to a real type

8594

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

8604–8611

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

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

8607

Create directly as of CGF.Int64Ty type.

8632

Expand auto here to a real type

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8580

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;
...

8634–8637

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;
...

8645–8646

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

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8594

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.

8604–8611

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

8607

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

8645–8646

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

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8107

Can we encapsulate Dims into StructNonContiguousInfo?

8604–8605

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

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

8604–8611

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

8607

Hmm, why, can you investigate?

8642–8645

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;
...

8645–8646

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?

8684

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

[cchen]

Resolve issues

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8607

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)

8642–8645

Not sure about this one, I've added:

if (!OASE)
  continue;
...

8645–8646

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

8607

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.

8645–8646

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.

8676

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

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8607

I'll investigate it, thanks.

8645–8646

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

8693

Use preincrement

9036

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

9690–9692

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

8615–8617

No need for braces here

9747

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

11129–11132

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

11373–11376

Same, pass the fields as arguments instead.

clang/lib/Serialization/ASTReader.cpp

12482

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

12599–12600

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

12656

.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);

12713

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

12717–12718

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

12756–12757

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

[cchen]

Fix coding style

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9573–9575

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?

9693–9721

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

16540

Better to use integer value as selectors, not boolean.

18399–18400

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

18455

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

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9573–9575

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

16540

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

16540

No, it is an integer, starts from 0

[cchen]

Fix based on feedback

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

11127–11128

Better to pass Info here directly.

[cchen]

Pass Info directly

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

16574

/*IsNonContiguous=*/false

16594

/*IsNonContiguous=*/false

18389–18390

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

[cchen]

Fix based on feedback

[cchen]

Rebase and resolve conflictions

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7823–7826

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

7823–7826

I rather doubt. You can try to ask @kkwli0

[kkwli0]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7823–7826

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

 //===- OpenMPClause.h - Classes for OpenMP clauses --------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 /// \file
class MappableComponent {
 
     /// 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),
           AssociatedDeclaration(
               AssociatedDeclaration
                   ? cast<ValueDecl>(AssociatedDeclaration->getCanonicalDecl())
                   : nullptr) {}
 
     Expr *getAssociatedExpression() const { return AssociatedExpression; }
 
     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.
   using MappableExprComponentList = SmallVector<MappableComponent, 8>;
   using MappableExprComponentListRef = ArrayRef<MappableComponent>;
 
   // List of all component lists associated to the same base declaration.
void setComponents(ArrayRef<MappableComponent> Components,
                      ArrayRef<unsigned> CLSs) {
     assert(Components.size() == NumComponents &&
            "Unexpected amount of component lists.");
     assert(CLSs.size() == NumComponentLists &&
            "Unexpected amount of list sizes.");
     std::copy(Components.begin(), Components.end(), getComponentsRef().begin());
   }
 
+  /// Get the non-contiguous attribute per declaration that are in the trailing
+  /// objects of the class.
+  MutableArrayRef<bool> getNonContiguousListsRef() {
+    return MutableArrayRef<bool>(
+        static_cast<T *>(this)->template getTrailingObjects<bool>(),
+        NumComponentLists);
+  }
+
+  /// Get the non-contiguous attribute per declaration that are in the trailing
+  /// objects of the class.
+  ArrayRef<bool> getNonContiguousListsRef() const {
+    return ArrayRef<bool>(
+        static_cast<const T *>(this)->template getTrailingObjects<bool>(),
+        NumComponentLists);
+  }
+
+  /// Set the non-contiguous attribute per declaration that are in the trailing
+  /// objects of the class.
+  void setNonContiguousLists(ArrayRef<bool> NLs) {
+    assert(NLs.size() == NumComponentLists &&
+           "Unexpected amount of list numbers.");
+    std::copy(NLs.begin(), NLs.end(), getNonContiguousListsRef().begin());
+  }
+
   /// Fill the clause information from the list of declarations and
   /// associated component lists.
   void setClauseInfo(ArrayRef<ValueDecl *> Declarations,
-                     MappableExprComponentListsRef ComponentLists) {
+                     MappableExprComponentListsRef ComponentLists,
+                     ArrayRef<bool> NonContiguousList) {
     // Perform some checks to make sure the data sizes are consistent with the
     // information available when the clause was created.
     assert(getUniqueDeclarationsTotalNumber(Declarations) ==
                NumUniqueDeclarations &&
            "Unexpected number of mappable expression info entries!");
     assert(getComponentsTotalNumber(ComponentLists) == NumComponents &&
            "Unexpected total number of components!");
     assert(Declarations.size() == ComponentLists.size() &&
            "Declaration and component lists size is not consistent!");
     assert(Declarations.size() == NumComponentLists &&
            "Unexpected declaration and component lists size!");
+    assert(NonContiguousList.size() == ComponentLists.size() &&
+           "Unexpected NonContiguousList size");
 
     // Organize the components by declaration and retrieve the original
     // expression. Original expressions are always the first component of the
     // mappable component list.
     llvm::MapVector<ValueDecl *, SmallVector<MappableExprComponentListRef, 8>>
         ComponentListMap;
     {
       auto CI = ComponentLists.begin();
for (auto &M : ComponentListMap) {
         // Save the size.
         *CLSI = PrevSize;
         ++CLSI;
 
         // Append components after the current components iterator.
         CI = std::copy(C.begin(), C.end(), CI);
       }
     }
+
+    std::copy(NonContiguousList.begin(), NonContiguousList.end(),
+              getNonContiguousListsRef().begin());
   }
 
   /// Set the nested name specifier of associated user-defined mapper.
   void setMapperQualifierLoc(NestedNameSpecifierLoc NNSL) {
     MapperQualifierLoc = NNSL;
   }
 
   /// Set the name of associated user-defined mapper.
public:
   using const_all_components_range =
       llvm::iterator_range<const_all_components_iterator>;
 
   const_all_components_range all_components() const {
     auto A = getComponentsRef();
     return const_all_components_range(A.begin(), A.end());
   }
 
+  using non_contiguous_list_iterator = MutableArrayRef<bool>::iterator;
+  using non_contiguous_list_const_iterator = ArrayRef<bool>::iterator;
+  using non_contiguous_list_range =
+      llvm::iterator_range<non_contiguous_list_iterator>;
+  using non_contiguous_list_const_range =
+      llvm::iterator_range<non_contiguous_list_const_iterator>;
+
+  non_contiguous_list_iterator non_contiguous_list_begin() {
+    return getNonContiguousListsRef().begin();
+  }
+  non_contiguous_list_iterator non_contiguous_list_end() {
+    return getNonContiguousListsRef().end();
+  }
+  non_contiguous_list_const_iterator non_contiguous_list_begin() const {
+    return getNonContiguousListsRef().begin();
+  }
+  non_contiguous_list_const_iterator non_contiguous_list_end() const {
+    return getNonContiguousListsRef().end();
+  }
+  non_contiguous_list_range non_contiguous_lists() {
+    return non_contiguous_list_range(non_contiguous_list_begin(),
+                                     non_contiguous_list_end());
+  }
+  non_contiguous_list_const_range non_contiguous_lists() const {
+    return non_contiguous_list_const_range(non_contiguous_list_begin(),
+                                           non_contiguous_list_end());
+  }
+
   using mapperlist_iterator = MutableArrayRef<Expr *>::iterator;
   using mapperlist_const_iterator = ArrayRef<const Expr *>::iterator;
   using mapperlist_range = llvm::iterator_range<mapperlist_iterator>;
   using mapperlist_const_range =
       llvm::iterator_range<mapperlist_const_iterator>;
 
   mapperlist_iterator mapperlist_begin() { return getUDMapperRefs().begin(); }
   mapperlist_iterator mapperlist_end() { return getUDMapperRefs().end(); }

 /// This represents clause 'map' in the '#pragma omp ...'
 /// directives.
 ///
 /// \code
 /// #pragma omp target map(a,b)
 /// \endcode
 /// 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> {
+class OMPMapClause final
+    : public OMPMappableExprListClause<OMPMapClause>,
+          OMPMapClause, Expr *, ValueDecl *, bool, unsigned,
   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();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
+  size_t numTrailingObjects(OverloadToken<bool>) const {
+    return getTotalComponentListNum();
+  }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 private:
   /// Map-type-modifiers for the 'map' clause.
   OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
public:
   /// \param MapperId The identifier of associated user-defined mapper.
   /// \param Type Map type.
   /// \param TypeIsImplicit Map type is inferred implicitly.
   /// \param TypeLoc Location of the map type.
   static OMPMapClause *
   Create(const ASTContext &C, const OMPVarListLocTy &Locs,
          ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
          MappableExprComponentListsRef ComponentLists,
-         ArrayRef<Expr *> UDMapperRefs,
+         ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
          ArrayRef<OpenMPMapModifierKind> MapModifiers,
          ArrayRef<SourceLocation> MapModifiersLoc,
          NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
          OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc);
 
   /// Creates an empty clause with the place for \a NumVars original
   /// expressions, \a NumUniqueDeclarations declarations, \NumComponentLists
   /// lists, and \a NumComponents expression components.

 ///
 /// \code
 /// #pragma omp target update to(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target update' has clause 'to'
 /// with the variables 'a' and 'b'.
 class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
                           private llvm::TrailingObjects<
-                              OMPToClause, Expr *, ValueDecl *, unsigned,
+                              OMPToClause, Expr *, ValueDecl *, bool, unsigned,
                               OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// Build clause with number of variables \a NumVars.
   ///
class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
   size_t numTrailingObjects(OverloadToken<Expr *>) const {
     // There are varlist_size() of expressions, and varlist_size() of
     // user-defined mappers.
     return 2 * varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
+  size_t numTrailingObjects(OverloadToken<bool>) const {
+    return getTotalComponentListNum();
+  }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// Creates clause with a list of variables \a Vars.
   ///
   /// \param C AST context.
   /// \param Locs Locations needed to build a mappable clause. It includes 1)
   /// StartLoc: starting location of the clause (the clause keyword); 2)
   /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
   /// \param Vars The original expression used in the clause.
   /// \param Declarations Declarations used in the clause.
   /// \param ComponentLists Component lists used in the clause.
   /// \param UDMapperRefs References to user-defined mappers associated with
   /// expressions used in the clause.
   /// \param UDMQualifierLoc C++ nested name specifier for the associated
   /// user-defined mapper.
   /// \param MapperId The identifier of associated user-defined mapper.
   static OMPToClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs,
                              ArrayRef<Expr *> Vars,
                              ArrayRef<ValueDecl *> Declarations,
                              MappableExprComponentListsRef ComponentLists,
+                             ArrayRef<bool> IsNonContiguousList,
                              ArrayRef<Expr *> UDMapperRefs,
                              NestedNameSpecifierLoc UDMQualifierLoc,
                              DeclarationNameInfo MapperId);
 
   /// Creates an empty clause with the place for \a NumVars variables.
   ///
   /// \param C AST context.
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)

 /// \code
 /// #pragma omp target update from(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target update' has clause 'from'
 /// with the variables 'a' and 'b'.
 class OMPFromClause final
     : public OMPMappableExprListClause<OMPFromClause>,
       private llvm::TrailingObjects<
-          OMPFromClause, Expr *, ValueDecl *, unsigned,
+          OMPFromClause, Expr *, ValueDecl *, bool, unsigned,
           OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// Build clause with number of variables \a NumVars.
   ///
class OMPFromClause final
   size_t numTrailingObjects(OverloadToken<Expr *>) const {
     // There are varlist_size() of expressions, and varlist_size() of
     // user-defined mappers.
     return 2 * varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
+  size_t numTrailingObjects(OverloadToken<bool>) const {
+    return getTotalComponentListNum();
+  }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// Creates clause with a list of variables \a Vars.
   ///
   /// \param C AST context.
   /// \param Locs Locations needed to build a mappable clause. It includes 1)
   /// StartLoc: starting location of the clause (the clause keyword); 2)
   /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
   /// \param Vars The original expression used in the clause.
   /// \param Declarations Declarations used in the clause.
   /// \param ComponentLists Component lists used in the clause.
   /// \param UDMapperRefs References to user-defined mappers associated with
   /// expressions used in the clause.
   /// \param UDMQualifierLoc C++ nested name specifier for the associated
   /// user-defined mapper.
   /// \param MapperId The identifier of associated user-defined mapper.
   static OMPFromClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs,
                                ArrayRef<Expr *> Vars,
                                ArrayRef<ValueDecl *> Declarations,
                                MappableExprComponentListsRef ComponentLists,
+                               ArrayRef<bool> NonContiguousList,
                                ArrayRef<Expr *> UDMapperRefs,
                                NestedNameSpecifierLoc UDMQualifierLoc,
                                DeclarationNameInfo MapperId);
 
   /// Creates an empty clause with the place for \a NumVars variables.
   ///
   /// \param C AST context.
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)

 /// \code
 /// #pragma omp target data use_device_ptr(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target data' has clause
 /// 'use_device_ptr' with the variables 'a' and 'b'.
 class OMPUseDevicePtrClause final
     : public OMPMappableExprListClause<OMPUseDevicePtrClause>,
       private llvm::TrailingObjects<
-          OMPUseDevicePtrClause, Expr *, ValueDecl *, unsigned,
+          OMPUseDevicePtrClause, Expr *, ValueDecl *, bool, unsigned,
           OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// Build clause with number of variables \a NumVars.
   ///
class OMPUseDevicePtrClause final
   /// 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 {
     return 3 * varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
+  size_t numTrailingObjects(OverloadToken<bool>) const {
+    return getTotalComponentListNum();
+  }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
   /// Sets the list of references to private copies with initializers for new
   /// private variables.
   /// \param VL List of references.
   void setPrivateCopies(ArrayRef<Expr *> VL);

 /// \code
 /// #pragma omp target is_device_ptr(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target' has clause
 /// 'is_device_ptr' with the variables 'a' and 'b'.
 class OMPIsDevicePtrClause final
     : public OMPMappableExprListClause<OMPIsDevicePtrClause>,
       private llvm::TrailingObjects<
-          OMPIsDevicePtrClause, Expr *, ValueDecl *, unsigned,
+          OMPIsDevicePtrClause, Expr *, ValueDecl *, bool, unsigned,
           OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// Build clause with number of variables \a NumVars.
   ///
class OMPIsDevicePtrClause final
   /// 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 {
     return varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
+  size_t numTrailingObjects(OverloadToken<bool>) const {
+    return getTotalComponentListNum();
+  }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// Creates clause with a list of variables \a Vars.
   ///
   /// \param C AST context.

clang/lib/AST/OpenMPClause.cpp

 //===- OpenMPClause.cpp - Classes for OpenMP clauses ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the subclesses of Stmt class declared in OpenMPClause.h
for (const ValueDecl *D : Declarations) {
     Cache.insert(VD);
   }
   return TotalNum;
 }
 
 OMPMapClause *OMPMapClause::Create(
     const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
     ArrayRef<ValueDecl *> Declarations,
-    MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+    MappableExprComponentListsRef ComponentLists,
+    ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
     ArrayRef<OpenMPMapModifierKind> MapModifiers,
     ArrayRef<SourceLocation> MapModifiersLoc,
     NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
     OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
   Sizes.NumComponentLists = ComponentLists.size();
   Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // 2 x NumVars x Expr* - we have an original list expression and an associated
   // user-defined mapper for each clause list entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
+  // NumComponentLists x bool - number of non-contiguous attribute.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   OMPMapClause *Clause = new (Mem)
       OMPMapClause(MapModifiers, MapModifiersLoc, UDMQualifierLoc, MapperId,
                    Type, TypeIsImplicit, TypeLoc, Locs, Sizes);
 
   Clause->setVarRefs(Vars);
   Clause->setUDMapperRefs(UDMapperRefs);
-  Clause->setClauseInfo(Declarations, ComponentLists);
+  Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
   Clause->setMapType(Type);
   Clause->setMapLoc(TypeLoc);
   return Clause;
 }
 
 OMPMapClause *
 OMPMapClause::CreateEmpty(const ASTContext &C,
                           const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumVars,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPMapClause(Sizes);
 }
 
 OMPToClause *OMPToClause::Create(
     const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
     ArrayRef<ValueDecl *> Declarations,
-    MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+    MappableExprComponentListsRef ComponentLists,
+    ArrayRef<bool> NonContiguousList, ArrayRef<Expr *> UDMapperRefs,
     NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
   Sizes.NumComponentLists = ComponentLists.size();
   Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // 2 x NumVars x Expr* - we have an original list expression and an associated
   // user-defined mapper for each clause list entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
+  // NumComponentLists x bool - number of non-contiguous attribute.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
 
   auto *Clause = new (Mem) OMPToClause(UDMQualifierLoc, MapperId, Locs, Sizes);
 
   Clause->setVarRefs(Vars);
   Clause->setUDMapperRefs(UDMapperRefs);
-  Clause->setClauseInfo(Declarations, ComponentLists);
+  Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
+
   return Clause;
 }
 
 OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C,
                                       const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPToClause(Sizes);
 }
 
 OMPFromClause *OMPFromClause::Create(
     const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
     ArrayRef<ValueDecl *> Declarations,
-    MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+    MappableExprComponentListsRef ComponentLists,
+    ArrayRef<bool> NonContiguousList,
+    ArrayRef<Expr *> UDMapperRefs,
     NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
   Sizes.NumComponentLists = ComponentLists.size();
   Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // 2 x NumVars x Expr* - we have an original list expression and an associated
   // user-defined mapper for each clause list entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
+  // NumComponentLists x bool - number of non-contiguous attribute.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
 
   auto *Clause =
       new (Mem) OMPFromClause(UDMQualifierLoc, MapperId, Locs, Sizes);
 
   Clause->setVarRefs(Vars);
   Clause->setUDMapperRefs(UDMapperRefs);
-  Clause->setClauseInfo(Declarations, ComponentLists);
+  Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
+
   return Clause;
 }
 
 OMPFromClause *
 OMPFromClause::CreateEmpty(const ASTContext &C,
                            const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPFromClause(Sizes);
 }
 
 void OMPUseDevicePtrClause::setPrivateCopies(ArrayRef<Expr *> VL) {
   assert(VL.size() == varlist_size() &&
          "Number of private copies is not the same as the preallocated buffer");
OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(
   Sizes.NumComponentLists = ComponentLists.size();
   Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // 3 x NumVars x Expr* - we have an original list expression for each clause
   // list entry and an equal number of private copies and inits.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
+  // NumComponentLists x bool - number of non-contiguous attribute.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
 
   OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause(Locs, Sizes);
 
   Clause->setVarRefs(Vars);
   Clause->setPrivateCopies(PrivateVars);
   Clause->setInits(Inits);
-  Clause->setClauseInfo(Declarations, ComponentLists);
+  SmallVector<bool, 4> NonContiguousList(Declarations.size(), false);
+  Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
   return Clause;
 }
 
 OMPUseDevicePtrClause *
 OMPUseDevicePtrClause::CreateEmpty(const ASTContext &C,
                                    const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPUseDevicePtrClause(Sizes);
 }
 
 OMPIsDevicePtrClause *
 OMPIsDevicePtrClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs,
                              ArrayRef<Expr *> Vars,
                              ArrayRef<ValueDecl *> Declarations,
                              MappableExprComponentListsRef ComponentLists) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
   Sizes.NumComponentLists = ComponentLists.size();
   Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // NumVars x Expr* - we have an original list expression for each clause list
   // entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
+  // NumComponentLists x bool - number of non-contiguous attribute.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
 
   OMPIsDevicePtrClause *Clause = new (Mem) OMPIsDevicePtrClause(Locs, Sizes);
 
   Clause->setVarRefs(Vars);
-  Clause->setClauseInfo(Declarations, ComponentLists);
+  SmallVector<bool, 4> NonContiguousList(Declarations.size(), false);
+  Clause->setClauseInfo(Declarations, ComponentLists, NonContiguousList);
   return Clause;
 }
 
 OMPIsDevicePtrClause *
 OMPIsDevicePtrClause::CreateEmpty(const ASTContext &C,
                                   const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
-      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+      totalSizeToAlloc<Expr *, ValueDecl *, bool, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumComponentLists,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPIsDevicePtrClause(Sizes);
 }
 
 OMPNontemporalClause *OMPNontemporalClause::Create(const ASTContext &C,
                                                    SourceLocation StartLoc,
                                                    SourceLocation LParenLoc,

clang/lib/CodeGen/CGOpenMPRuntime.h

public:
     /// The array of base pointer passed to the runtime library.
     llvm::Value *BasePointersArray = nullptr;
     /// The array of section pointers passed to the runtime library.
     llvm::Value *PointersArray = nullptr;
     /// The array of sizes passed to the runtime library.
     llvm::Value *SizesArray = nullptr;
     /// The array of map types passed to the runtime library.
     llvm::Value *MapTypesArray = nullptr;
+    /// The array of array of dims passed to the runtime library.
+    llvm::Value *DimsArray = nullptr;
+    /// The array of array of descriptor passed to the runtime library.
+    llvm::Value *DescriptorsArray = 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;
 
     explicit TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo)
         : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}
     /// Clear information about the data arrays.
     void clearArrayInfo() {
       BasePointersArray = nullptr;
       PointersArray = nullptr;
       SizesArray = nullptr;
       MapTypesArray = nullptr;
+      DimsArray = nullptr;
+      DescriptorsArray = nullptr;
       NumberOfPtrs = 0u;
     }
     /// Return true if the current target data information has valid arrays.
     bool isValid() {
       return BasePointersArray && PointersArray && SizesArray &&
              MapTypesArray && NumberOfPtrs;
     }
     bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }

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 = 0x800,
     /// 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()};
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> HighestElem = {
         0, Address::invalid()};
     Address Base = Address::invalid();
   };
 
 private:
   /// Kind that defines how a device pointer has to be returned.
   struct MapInfo {
     OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
     ArrayRef<OpenMPMapModifierKind> MapModifiers;
     bool ReturnDevicePointer = false;
     bool IsImplicit = false;
+    bool IsNonContiguous = false;
 
     MapInfo() = default;
     MapInfo(
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
         OpenMPMapClauseKind MapType,
-        ArrayRef<OpenMPMapModifierKind> MapModifiers,
-        bool ReturnDevicePointer, bool IsImplicit)
+        ArrayRef<OpenMPMapModifierKind> MapModifiers, bool ReturnDevicePointer,
+        bool IsImplicit, bool IsNonContiguous)
         : Components(Components), MapType(MapType), MapModifiers(MapModifiers),
-          ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
+          ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
+          IsNonContiguous(IsNonContiguous) {}
   };
 
   /// 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 {
     const Expr *IE = nullptr;
     const ValueDecl *VD = nullptr;
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,
       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) 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;
     //
     // struct S1 {
     //   int i;
     //   float f[50];
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 =
-          isFinalArraySectionExpression(I->getAssociatedExpression());
+          isFinalArraySectionExpression(I->getAssociatedExpression()) &&
+          (!IsNonContiguous);
 
       // 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;
 
+      if (OASE || OAShE ||
+          dyn_cast<ArraySubscriptExpr>(I->getAssociatedExpression())) {
+        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()) ||
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=*/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,
+          OpenMPOffloadMappingFlags Flags =
+              getMapTypeBits(MapType, MapModifiers, IsImplicit,
                              !IsExpressionFirstInfo || RequiresReference,
-              IsCaptureFirstInfo && !RequiresReference);
+                             IsCaptureFirstInfo && !RequiresReference,
+                             /*IsNonContiguous=*/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);
 
for (; I != CE; ++I) {
           BP = LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
       }
     }
   }
 
+  /// Generate the base pointers, section pointers, sizes , map type bits,
+  /// dimension size, offset, count, and strides 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 generateInfoForTargetDataComponentList(
+      OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
+      OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
+      MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
+      MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, MapDimArrayTy &Dims,
+      MapNonContiguousArrayTy &Offsets, MapNonContiguousArrayTy &Counts,
+      MapNonContiguousArrayTy &Strides, StructRangeInfoTy &PartialStruct,
+      bool IsFirstComponentList, bool IsImplicit,
+      ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
+          OverlappedElements = llvm::None) const {
+
+    generateInfoForComponentList(MapType, MapModifiers, Components,
+                                 BasePointers, Pointers, Sizes, Types, Dims,
+                                 PartialStruct, IsFirstComponentList,
+                                 IsImplicit, OverlappedElements, true);
+
+    const ASTContext &Context = CGF.getContext();
+
+    MapValuesArrayTy CurOffsets;
+    MapValuesArrayTy CurCounts;
+    MapValuesArrayTy CurStrides;
+    llvm::Value *CurStride = nullptr;
+
+    // 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.
+    SmallVector<llvm::Value *, 4> DimSizes;
+    for (const auto &Component : Components) {
+      const Expr *AssocExpr = Component.getAssociatedExpression();
+      const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+      const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+      if (AE || OASE) {
+        QualType Ty;
+        if (OASE)
+          Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+        else
+          Ty = AE->getType();
+        auto *CAT = Context.getAsConstantArrayType(Ty);
+        auto *VAT = Context.getAsVariableArrayType(Ty);
+        // Get element size if we 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 {
+            assert(VAT && "Should be either ConstantArray or VariableArray");
+            ElementType = VAT->getElementType().getTypePtr();
+            ElementTypeSize =
+              Context.getTypeSizeInChars(ElementType).getQuantity();
+          }
+          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 {
+          assert(VAT && "Should be either ConstantArray or VariableArray");
+          const Expr *Size = VAT->getSizeExpr();
+          SizeV = CGF.EmitScalarExpr(Size);
+        }
+        SizeV = CGF.Builder.CreateIntCast(SizeV, CGF.Int64Ty,
+                                          /*IsSigned=*/false);
+        DimSizes.push_back(SizeV);
+      }
+    }
+
+    // Scan the components from the base to the complete expression.
+    auto CI = Components.begin();
+    auto CE = Components.end();
+    auto I = CI;
+    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 (; I != CE; ++I) {
+      const Expr *AssocExpr = I->getAssociatedExpression();
+      const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+      const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+      if (OASE || AE) {
+        // Offset
+        const Expr *OffsetExpr = nullptr;
+        if (OASE)
+          OffsetExpr = OASE->getLowerBound();
+        else
+          OffsetExpr = AE->getIdx();
+        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.begin()) {
+          CurStride =
+              CGF.Builder.CreateNUWMul(CurStrides.back(), *std::prev(DI, 1));
+          CurStrides.push_back(CurStride);
+        }
+
+        DI++;
+      }
+    }
+
+    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
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](
         const ValueDecl *D,
         OMPClauseMappableExprCommon::MappableExprComponentListRef L,
         OpenMPMapClauseKind MapType,
         ArrayRef<OpenMPMapModifierKind> MapModifiers,
-        bool ReturnDevicePointer, bool IsImplicit) {
+        bool ReturnDevicePointer, bool IsImplicit,
+        bool IsNonContiguous) {
       const ValueDecl *VD =
           D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
       Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
-                            IsImplicit);
+                            IsImplicit, IsNonContiguous);
     };
 
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
       for (const auto L : C->component_lists()) {
         InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
-            /*ReturnDevicePointer=*/false, C->isImplicit());
+                /*ReturnDevicePointer=*/false, C->isImplicit(),
+                /*IsNonContiguous=*/false);
       }
-    for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
-      for (const auto L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
-            /*ReturnDevicePointer=*/false, C->isImplicit());
+    for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>()) {
+      auto CI = C->component_lists_begin();
+      auto CE = C->component_lists_end();
+      auto NI = C->non_contiguous_list_begin();
+      for (; CI != CE; ++CI, ++NI) {
+        InfoGen((*CI).first, (*CI).second, OMPC_MAP_to, llvm::None,
+                /*ReturnDevicePointer=*/false, C->isImplicit(), *NI);
+      }
+    }
+    for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>()) {
+      auto CI = C->component_lists_begin();
+      auto CE = C->component_lists_end();
+      auto NI = C->non_contiguous_list_begin();
+      for (; CI != CE; ++CI, ++NI) {
+        InfoGen((*CI).first, (*CI).second, OMPC_MAP_from, llvm::None,
+                /*ReturnDevicePointer=*/false, C->isImplicit(), *NI);
       }
-    for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
-      for (const auto L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
-            /*ReturnDevicePointer=*/false, C->isImplicit());
-      }
 
     // Look at the use_device_ptr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_ptr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
     // emission of that entry until the whole struct has been processed.
     llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
for (const auto *C :
         // 'this', which maps to null in the map of map information. Otherwise
         // look directly for the information.
         auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
 
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
           auto CI = std::find_if(
-              It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {
+              It->second.begin(), It->second.end(),
+              [VD](const MapInfo &MI) {
                 return MI.Components.back().getAssociatedDeclaration() == VD;
               });
           // If we found a map entry, signal that the pointer has to be returned
           // and move on to the next declaration.
           if (CI != It->second.end()) {
             CI->ReturnDevicePointer = true;
             continue;
           }
         }
 
         // We didn't find any match in our map information - generate a zero
         // size array section - if the pointer is a struct member we defer this
         // action until the whole struct has been processed.
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
           InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None,
-                  /*ReturnDevicePointer=*/false, C->isImplicit());
+                  /*ReturnDevicePointer=*/false, C->isImplicit(),
+                  /*IsNonContiguous=*/false);
           DeferredInfo[nullptr].emplace_back(IE, VD);
         } else {
           llvm::Value *Ptr =
               CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
           BasePointers.emplace_back(Ptr, VD);
           Pointers.push_back(Ptr);
           Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
           Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM);
         }
       }
     }
 
     for (const auto &M : Info) {
       // We need to know when we generate information for the first component
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // 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();
+        if (L.IsNonContiguous) {
+          generateInfoForTargetDataComponentList(
+              L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+              CurPointers, CurSizes, CurTypes, CurDims, CurOffsets, CurCounts,
+              CurStrides, PartialStruct, IsFirstComponentList, L.IsImplicit);
+        } else {
+          // Indicate that we do not do the special non-contiguous codegen
           generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
                                        CurBasePointers, CurPointers, CurSizes,
-                                     CurTypes, PartialStruct,
+                                       CurTypes, CurDims, PartialStruct,
                                        IsFirstComponentList, L.IsImplicit);
+        }
 
         // 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();
for (const auto &M : Info) {
           // correct value of MEMBER_OF.
           CurTypes.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM |
                              OMP_MAP_MEMBER_OF);
         }
       }
 
       // 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 `emitTargetDataOffloadingArrays` use it
+        // to skip this one.
+        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());
+      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,
auto &&InfoGen = [&Info](
         const ValueDecl *D,
         OMPClauseMappableExprCommon::MappableExprComponentListRef L,
         OpenMPMapClauseKind MapType,
         ArrayRef<OpenMPMapModifierKind> MapModifiers,
         bool ReturnDevicePointer, bool IsImplicit) {
       const ValueDecl *VD =
           D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
       Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
-                            IsImplicit);
+                            IsImplicit, /*IsNonContiguous=*/false);
     };
 
     for (const auto *C : CurMapperDir->clauselists()) {
       const auto *MC = cast<OMPMapClause>(C);
       for (const auto L : MC->component_lists()) {
         InfoGen(L.first, L.second, MC->getMapType(), MC->getMapTypeModifiers(),
                 /*ReturnDevicePointer=*/false, MC->isImplicit());
       }
     }
 
     for (const auto &M : Info) {
       // We need to know when we generate information for the first component
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // Temporary versions of arrays
       MapBaseValuesArrayTy CurBasePointers;
       MapValuesArrayTy CurPointers;
       MapValuesArrayTy CurSizes;
       MapFlagsArrayTy CurTypes;
+      MapDimArrayTy CurDims;
       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())
         emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
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) {
         Address SAddr(S, Ctx.getTypeAlignInChars(Int64Ty));
         CGF.Builder.CreateStore(
             CGF.Builder.CreateIntCast(Sizes[I], CGM.Int64Ty, /*isSigned=*/true),
             SAddr);
       }
     }
   }
 }
 
+/// 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
+emitTargetDataOffloadingArrays(CodeGenFunction &CGF,
+                     MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
+                     MappableExprsHandler::MapValuesArrayTy &Pointers,
+                     MappableExprsHandler::MapValuesArrayTy &Sizes,
+                     MappableExprsHandler::MapFlagsArrayTy &MapTypes,
+                     MappableExprsHandler::MapDimArrayTy &Dims,
+                     MappableExprsHandler::MapNonContiguousArrayTy &Offsets,
+                     MappableExprsHandler::MapNonContiguousArrayTy &Counts,
+                     MappableExprsHandler::MapNonContiguousArrayTy &Strides,
+                     CGOpenMPRuntime::TargetDataInfo &Info) {
+  emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Dims, Info,
+                       true);
+
+  if (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.NumberOfPtrs; 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(Offsets[L][RevIdx], OffsetLVal);
+        // Count
+        LValue CountLVal = CGF.EmitLValueForField(
+            DimsLVal, *std::next(RD->field_begin(), CountFD));
+        CGF.EmitStoreOfScalar(Counts[L][RevIdx], CountLVal);
+        // Stride
+        LValue StrideLVal = CGF.EmitLValueForField(
+            DimsLVal, *std::next(RD->field_begin(), StrideFD));
+        CGF.EmitStoreOfScalar(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) {
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 the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitTargetDataOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes,
+                                   Dims, Offsets, Counts, Strides, Info);
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
                                  SizesArrayArg, MapTypesArrayArg, Info);
 
     // Emit device ID if any.
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 the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitTargetDataOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes,
+                                   Dims, Offsets, Counts, Strides, 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());
     InputInfo.SizesArray =
         Address(Info.SizesArray, 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);
+    bool &IsNonContiguous, 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())) {
         // TODO: try to make it firstprivate.
         if (DVar.CKind != OMPC_unknown)
           ImplicitFirstprivate.push_back(E);
       }
       return;
     }
     if (isOpenMPTargetExecutionDirective(DKind)) {
       OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
-      if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+      bool IsNonContiguous = false;
+      if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents,
+                                        IsNonContiguous, 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 &IsNonContiguousRef;
   bool NoDiagnose = false;
   const Expr *RelevantExpr = nullptr;
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
   SourceLocation ELoc;
   SourceRange ERange;
 
   void emitErrorMsg() {
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;
     }
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
       Expr::EvalResult ResultR;
       Expr::EvalResult ResultL;
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);
     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);
 
     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);
     }
     return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
   }
   bool VisitBinaryOperator(BinaryOperator *BO) {
     if (SemaRef.getLangOpts().OpenMP < 50 || !BO->getType()->isPointerType()) {
       emitErrorMsg();
       return false;
     }
public:
   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),
-        NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
+      bool &IsNonContiguousTargetUpdate, bool NoDiagnose, SourceLocation &ELoc,
+      SourceRange &ERange)
+      : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components),
+        IsNonContiguousRef(IsNonContiguousTargetUpdate), 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) {
+    bool &IsNonContiguousTargetUpdate, OpenMPClauseKind CKind,
+    OpenMPDirectiveKind DKind, bool NoDiagnose) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
-  MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
-                         ERange);
+  MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents,
+                         IsNonContiguousTargetUpdate, 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.
 static bool checkMapConflicts(
struct MappableVarListInfo {
   // The list of processed expressions.
   SmallVector<Expr *, 16> ProcessedVarList;
   // The mappble components for each expression.
   OMPClauseMappableExprCommon::MappableExprComponentLists VarComponents;
   // The base declaration of the variable.
   SmallVector<ValueDecl *, 16> VarBaseDeclarations;
   // The reference to the user-defined mapper associated with every expression.
   SmallVector<Expr *, 16> UDMapperList;
+  // The list of whether the expression is non-contiguous or not
+  SmallVector<bool, 16> IsNonContiguousList;
 
   MappableVarListInfo(ArrayRef<Expr *> VarList) : VarList(VarList) {
     // We have a list of components and base declarations for each entry in the
     // variable list.
     VarComponents.reserve(VarList.size());
     VarBaseDeclarations.reserve(VarList.size());
   }
 };
if (!RE->isLValue()) {
       } else {
         SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses)
             << getOpenMPClauseName(CKind) << RE->getSourceRange();
       }
       continue;
     }
 
     OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
+    bool IsNonContiguousTargetUpdate = false;
     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, IsNonContiguousTargetUpdate, 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
       DSAS->addMappedClassesQualTypes(TE->getType());
       // Try to find the associated user-defined mapper.
       ExprResult ER = buildUserDefinedMapperRef(
           SemaRef, DSAS->getCurScope(), MapperIdScopeSpec, MapperId,
           VE->getType().getCanonicalType(), UnresolvedMapper);
       if (ER.isInvalid())
         continue;
       MVLI.UDMapperList.push_back(ER.get());
       // Skip restriction checking for variable or field declarations
       MVLI.ProcessedVarList.push_back(RE);
       MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
       MVLI.VarComponents.back().append(CurComponents.begin(),
                                        CurComponents.end());
       MVLI.VarBaseDeclarations.push_back(nullptr);
+      MVLI.IsNonContiguousList.push_back(IsNonContiguousTargetUpdate);
       continue;
     }
 
     // For the following checks, we rely on the base declaration which is
     // expected to be associated with the last component. The declaration is
     // expected to be a variable or a field (if 'this' is being mapped).
     CurDeclaration = CurComponents.back().getAssociatedDeclaration();
     assert(CurDeclaration && "Null decl on map clause.");
for (Expr *RE : MVLI.VarList) {
     // Save the components and declaration to create the clause. For purposes of
     // the clause creation, any component list that has has base 'this' uses
     // null as base declaration.
     MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
     MVLI.VarComponents.back().append(CurComponents.begin(),
                                      CurComponents.end());
     MVLI.VarBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr
                                                            : CurDeclaration);
+    MVLI.IsNonContiguousList.push_back(IsNonContiguousTargetUpdate);
   }
 }
 
 OMPClause *Sema::ActOnOpenMPMapClause(
     ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
     ArrayRef<SourceLocation> MapTypeModifiersLoc,
     CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
OMPClause *Sema::ActOnOpenMPMapClause(
 
   MappableVarListInfo MVLI(VarList);
   checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc,
                               MapperIdScopeSpec, MapperId, UnresolvedMappers,
                               MapType, IsMapTypeImplicit);
 
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
-  return OMPMapClause::Create(Context, Locs, MVLI.ProcessedVarList,
-                              MVLI.VarBaseDeclarations, MVLI.VarComponents,
-                              MVLI.UDMapperList, Modifiers, ModifiersLoc,
-                              MapperIdScopeSpec.getWithLocInContext(Context),
-                              MapperId, MapType, IsMapTypeImplicit, MapLoc);
+  return OMPMapClause::Create(
+      Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
+      MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
+      Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context),
 }
 
 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
                                                TypeResult ParsedType) {
   assert(ParsedType.isUsable());
 
   QualType ReductionType = GetTypeFromParser(ParsedType.get());
   if (ReductionType.isNull())
OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
   MappableVarListInfo MVLI(VarList);
   checkMappableExpressionList(*this, DSAStack, OMPC_to, MVLI, Locs.StartLoc,
                               MapperIdScopeSpec, MapperId, UnresolvedMappers);
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPToClause::Create(
       Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
-      MVLI.VarComponents, MVLI.UDMapperList,
+      MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
       MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
 }
 
 OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
                                        CXXScopeSpec &MapperIdScopeSpec,
                                        DeclarationNameInfo &MapperId,
                                        const OMPVarListLocTy &Locs,
                                        ArrayRef<Expr *> UnresolvedMappers) {
   MappableVarListInfo MVLI(VarList);
   checkMappableExpressionList(*this, DSAStack, OMPC_from, MVLI, Locs.StartLoc,
                               MapperIdScopeSpec, MapperId, UnresolvedMappers);
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPFromClause::Create(
       Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
-      MVLI.VarComponents, MVLI.UDMapperList,
+      MVLI.VarComponents, MVLI.IsNonContiguousList, MVLI.UDMapperList,
       MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
 }
 
 OMPClause *Sema::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
                                                const OMPVarListLocTy &Locs) {
   MappableVarListInfo MVLI(VarList);
   SmallVector<Expr *, 8> PrivateCopies;
   SmallVector<Expr *, 8> Inits;
for (Expr *RefExpr : VarList) {
     // 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));
   }
 
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDevicePtrClause::Create(
       Context, Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
       MVLI.VarBaseDeclarations, MVLI.VarComponents);
 }
 
 OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                               const OMPVarListLocTy &Locs) {
   MappableVarListInfo MVLI(VarList);
   for (Expr *RefExpr : VarList) {
     assert(RefExpr && "NULL expr in OpenMP is_device_ptr clause.");
     SourceLocation ELoc;
     SourceRange ERange;
     Expr *SimpleRefExpr = RefExpr;
     auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange);
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);
     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

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();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
     Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
         AssociatedExpr, AssociatedDecl));
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPAllocateClause(OMPAllocateClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   C->setColonLoc(Record.readSourceLocation());
   C->setAllocator(Record.readSubExpr());
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));
   }
   C->setComponents(Components, ListSizes);
+
+  SmallVector<bool, 16> ListNonContiguous;
+  ListNonContiguous.reserve(TotalLists);
+  for (unsigned i = 0; i < TotalLists; ++i) {
+    ListNonContiguous.push_back(Record.readBool());
+  }
+  C->setNonContiguousLists(ListNonContiguous);
 }
 
 void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
   C->setMapperIdInfo(Record.readDeclarationNameInfo());
   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));
   }
   C->setComponents(Components, ListSizes);
+
+  SmallVector<bool, 16> ListNonContiguous;
+  ListNonContiguous.reserve(TotalLists);
+  for (unsigned i = 0; i < TotalLists; ++i) {
+    ListNonContiguous.push_back(Record.readBool());
+  }
+  C->setNonContiguousLists(ListNonContiguous);
 }
 
 void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
   auto TotalLists = C->getTotalComponentListNum();
   auto TotalComponents = C->getTotalComponentsNum();
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));
   }
   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();
   auto TotalComponents = C->getTotalComponentsNum();
 
   SmallVector<Expr *, 16> Vars;
   Vars.reserve(NumVars);
void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
   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));
   }
   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.AddDeclRef(M.getAssociatedDeclaration());
   }
+  for (auto NC : C->non_contiguous_lists())
+    Record.push_back(NC);
 }
 
 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.AddDeclRef(M.getAssociatedDeclaration());
   }
+  for (auto NC : C->non_contiguous_lists())
+    Record.push_back(NC);
 }
 
 void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *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());

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 2081]
+
+// 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 2081]
+
+// 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
 #endif