[OpenMP 5.0] Fix user-defined mapper lookup in sema

[lildmh]

Summary

This patches fixes the case when a user-defined mapper is attached to the elements of an array, and to report error when a mapper is used for types other than struct, class, and union.

[ABataev]

What happens without this patch?

[lildmh]

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

[ABataev]

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

[lildmh]

In D67978#1681319, @ABataev wrote:

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

Without this patch, it won't be able to find the mapper in this case, and Clang assumes no mapper. This patch fixes it, so the compiler can find the mapper in this case.

[ABataev]

In D67978#1681333, @lildmh wrote:

In D67978#1681319, @ABataev wrote:

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

Without this patch, it won't be able to find the mapper in this case, and Clang assumes no mapper. This patch fixes it, so the compiler can find the mapper in this case.

It means, that you just ignore mappers for the array types, right?

[lildmh]

In D67978#1681388, @ABataev wrote:

In D67978#1681333, @lildmh wrote:

In D67978#1681319, @ABataev wrote:

In D67978#1681318, @lildmh wrote:

Without this patch, it cannot recognize array with mapper, for instance, #pragma omp target map(mapper(a),to: arr[0:2]) won't work without this patch.

What if we have a mapper for the array type?

Without this patch, it won't be able to find the mapper in this case, and Clang assumes no mapper. This patch fixes it, so the compiler can find the mapper in this case.

It means, that you just ignore mappers for the array types, right?

Yes, it ignored mappers before, and this patch fixes it.

[ABataev]

Definitely need positive tests with ast printing and codegen.

lib/Sema/SemaOpenMP.cpp

14751 ↗

(On Diff #221577)

Why do you want canonical type here? I think it is wrong. It drops all language sugar like typedefs etc. But typedefs are not supported in mappers, right?

[lildmh]

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

lib/Sema/SemaOpenMP.cpp

14751 ↗

(On Diff #221577)

I didn't see that the spec says typedef is not supported in mappers, so I suppose it should be supported. So I think getCanonicalType is necessary here?

[ABataev]

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

lib/Sema/SemaOpenMP.cpp

14751 ↗

(On Diff #221577)

Do we check the canonicaL type by default, for non-array type?

[lildmh]

In D67978#1683166, @ABataev wrote:

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

The ast print for array with mapper was correct because the mapper id is still with the array type. Without this patch, the problem is it will not look up the mapper declaration associated with the id, and as a result, the codegen is not correct. I found this problem when I tested the codegen.

lib/Sema/SemaOpenMP.cpp

14751 ↗

(On Diff #221577)

Yes, we use canonical type for non-array type, for instance, line 14935 in this file.

[ABataev]

In D67978#1684104, @lildmh wrote:

In D67978#1683166, @ABataev wrote:

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

The ast print for array with mapper was correct because the mapper id is still with the array type. Without this patch, the problem is it will not look up the mapper declaration associated with the id, and as a result, the codegen is not correct. I found this problem when I tested the codegen.

Then another one question. Why we don't emit the diagnostics if the original type is not a class, struct or union? We just ignore this situation currently but we should not. The error message must be emitted. I think that's why the ast-print test works correctly though it should not.

[lildmh]

In D67978#1684169, @ABataev wrote:

In D67978#1684104, @lildmh wrote:

In D67978#1683166, @ABataev wrote:

In D67978#1683146, @lildmh wrote:

HI Alexey, the ast print test is already there. Because I didn't check the mapper for array type before, such code will always not report any error, and ast print test is correct. Codegen test belongs to the other patch I released. It fits that patch much better.

How is this possible? If we did not have support for the array type, we could not have correct handling of such types in successful tests.

The ast print for array with mapper was correct because the mapper id is still with the array type. Without this patch, the problem is it will not look up the mapper declaration associated with the id, and as a result, the codegen is not correct. I found this problem when I tested the codegen.

Then another one question. Why we don't emit the diagnostics if the original type is not a class, struct or union? We just ignore this situation currently but we should not. The error message must be emitted. I think that's why the ast-print test works correctly though it should not.

We emit such diagnostic when mapper is declared, i.e., you cannot define a mapper for a non-class, non-struct, non-union type. It's in function ActOnOpenMPDeclareMapperType in SemaOpenMP.cpp. But I didn't emit such information where mapper is used. I will fix it in this patch. Thanks for getting this!

[lildmh]

Fix mapper type checking

[ABataev]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

Why need an additional check for scope and not "default" id? I don't see this additional requirement in the standard.

[lildmh]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

It's because every variable in map clauses will check this, including those are not struct, class, or union.

Using this, e.g., mapping an integer won't report error that it doesn't have a mapper.

[ABataev]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

Maybe just move the check to the end of the function?

[lildmh]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

It will do some additional useless work if I move it to the end. I don't think it is necessary to move it back.

But if you believe it's better, I can do it.

[ABataev]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

Could you add tests for this new functionality?

[lildmh]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

I already add the test in declare_mapper_message.c. I think nothing is needed for ast print test.

[ABataev]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

I mean, for this error message. I did not see the change in the tests when you added this new check.

[lildmh]

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

line 55, 68, 79 in declare_mapper_messages.c is added for this new check.

[ABataev]

LG.

lib/Sema/SemaOpenMP.cpp

14805 ↗

(On Diff #221991)

Ah, ok, found it, thanks!

[lildmh]

Thanks Alexey! Please check the other 2 mapper patches at https://reviews.llvm.org/D67833 and https://reviews.llvm.org/D68100 when you have time. They should be the last mapper patches.

[ABataev]

In D67978#1684845, @lildmh wrote:

Thanks Alexey! Please check the other 2 mapper patches at https://reviews.llvm.org/D67833 and https://reviews.llvm.org/D68100 when you have time. They should be the last mapper patches.

Sure, will do this tomorrow.

cfe/trunk/lib/Sema/SemaOpenMP.cpp

 // build a reference to it.
 static ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S,
                                             CXXScopeSpec &MapperIdScopeSpec,
                                             const DeclarationNameInfo &MapperId,
                                             QualType Type,
                                             Expr *UnresolvedMapper) {
   if (MapperIdScopeSpec.isInvalid())
     return ExprError();
+  // Get the actual type for the array type.
+  if (Type->isArrayType()) {
+    assert(Type->getAsArrayTypeUnsafe() && "Expect to get a valid array type");
+    Type = Type->getAsArrayTypeUnsafe()->getElementType().getCanonicalType();
+  }
   // Find all user-defined mappers with the given MapperId.
   SmallVector<UnresolvedSet<8>, 4> Lookups;
   LookupResult Lookup(SemaRef, MapperId, Sema::LookupOMPMapperName);
   Lookup.suppressDiagnostics();
   if (S) {
     while (S && SemaRef.LookupParsedName(Lookup, S, &MapperIdScopeSpec)) {
       NamedDecl *D = Lookup.getRepresentativeDecl();
       while (S && !S->isDeclScope(D))
for (const UnresolvedSet<8> &Set : Lookups) {
         continue;
       URS.append(Set.begin(), Set.end());
     }
     return UnresolvedLookupExpr::Create(
         SemaRef.Context, /*NamingClass=*/nullptr,
         MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperId,
         /*ADL=*/false, /*Overloaded=*/true, URS.begin(), URS.end());
   }
+  SourceLocation Loc = MapperId.getLoc();
   // [OpenMP 5.0], 2.19.7.3 declare mapper Directive, Restrictions
   //  The type must be of struct, union or class type in C and C++
-  if (!Type->isStructureOrClassType() && !Type->isUnionType())
-    return ExprEmpty();
-  SourceLocation Loc = MapperId.getLoc();
+  if (!Type->isStructureOrClassType() && !Type->isUnionType() &&
+      (MapperIdScopeSpec.isSet() || MapperId.getAsString() != "default")) {
+    SemaRef.Diag(Loc, diag::err_omp_mapper_wrong_type);
+    return ExprError();
+  }
   // Perform argument dependent lookup.
   if (SemaRef.getLangOpts().CPlusPlus && !MapperIdScopeSpec.isSet())
     argumentDependentLookup(SemaRef, MapperId, Loc, Type, Lookups);
   // Return the first user-defined mapper with the desired type.
   if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
           Lookups, [&SemaRef, Type](ValueDecl *D) -> ValueDecl * {
             if (!D->isInvalidDecl() &&
                 SemaRef.Context.hasSameType(D->getType(), Type))

cfe/trunk/test/OpenMP/declare_mapper_messages.c

 int fun(int arg) {
 #pragma omp declare mapper(id: struct vec v) map(v.len)
   {
 #pragma omp declare mapper(id: struct vec v) map(v.len)                 // expected-note {{previous definition is here}}
 #pragma omp declare mapper(id: struct vec v) map(v.len)                 // expected-error {{redefinition of user-defined mapper for type 'struct vec' with name 'id'}}
     {
 #pragma omp declare mapper(id: struct vec v) map(v.len) allocate(v)   // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp declare mapper'}}
       struct vec vv, v1;
+      struct vec arr[10];
+      double d;
 #pragma omp target map(mapper)                                          // expected-error {{use of undeclared identifier 'mapper'}}
       {}
 #pragma omp target map(mapper:vv)                                       // expected-error {{expected '(' after 'mapper'}}
       {}
 #pragma omp target map(mapper( :vv)                                     // expected-error {{expected expression}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}}
       {}
 #pragma omp target map(mapper(aa :vv)                                   // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}}
       {}
 #pragma omp target map(mapper(ab) :vv)                                  // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}}
       {}
+#pragma omp target map(mapper(ab) :arr[0:2])                            // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}}
+      {}
 #pragma omp target map(mapper(aa) :vv)                                  // expected-error {{missing map type}}
       {}
-#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1)
+#pragma omp target map(mapper(aa) to:d)                                 // expected-error {{mapper type must be of struct, union or class type}}
+      {}
+#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) map(mapper(aa) to:arr[0])
       {}
 
 #pragma omp target update to(mapper)                                    // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper()                                   // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper:vv)                                 // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(:vv)                                // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(aa :vv)                             // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(ab):vv)                             // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-#pragma omp target update to(mapper(aa):vv)
+#pragma omp target update to(mapper(ab):arr[0:2])                       // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(mapper(aa) a:vv)                           // expected-warning {{missing ':' after ) - ignoring}}
+#pragma omp target update to(mapper(aa):d)                              // expected-error {{mapper type must be of struct, union or class type}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(mapper(aa):vv) to(mapper(aa):arr[0])
 
 #pragma omp target update from(mapper)                                  // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper()                                 // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper:vv)                               // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(:vv)                              // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(aa :vv)                           // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(ab):vv)                           // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(mapper(ab):arr[0:2])                     // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(aa) a:vv)                         // expected-warning {{missing ':' after ) - ignoring}}
-#pragma omp target update from(mapper(aa):vv)
+#pragma omp target update from(mapper(aa):d)                            // expected-error {{mapper type must be of struct, union or class type}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(mapper(aa):vv) from(mapper(aa):arr[0])
     }
   }
   return arg;
 }

cfe/trunk/test/OpenMP/declare_mapper_messages.cpp

 
 int fun(int arg) {
 #pragma omp declare mapper(id: vec v) map(v.len)
   {
 #pragma omp declare mapper(id: vec v) map(v.len)                        // expected-note {{previous definition is here}}
     {
 #pragma omp declare mapper(id: vec v) map(v.len)
       vec vv, v1;
+      vec arr[10];
 #pragma omp target map(mapper)                                          // expected-error {{use of undeclared identifier 'mapper'}}
       {}
 #pragma omp target map(mapper:vv)                                       // expected-error {{expected '(' after 'mapper'}}
       {}
 #pragma omp target map(mapper( :vv)                                     // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
       {}
 #pragma omp target map(mapper(aa :vv)                                   // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
       {}
 #pragma omp target map(mapper(ab) :vv)                                  // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}}
       {}
 #pragma omp target map(mapper(N2::) :vv)                                // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}}
       {}
 #pragma omp target map(mapper(N1::) :vv)                                // expected-error {{illegal OpenMP user-defined mapper identifier}}
       {}
 #pragma omp target map(mapper(aa) :vv)                                  // expected-error {{missing map type}}
       {}
 #pragma omp target map(mapper(N1::aa) alloc:vv)                         // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}}
       {}
-#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1)
+#pragma omp target map(mapper(N1::aa) alloc:arr[0:2])                   // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}}
+      {}
+#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) map(mapper(aa) to:arr[0])
       {}
 #pragma omp target map(mapper(N1::stack<int>::id) to:vv)
       {}
 
 #pragma omp target update to(mapper)                                    // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper()                                   // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper:vv)                                 // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(:vv)                                // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(aa :vv)                             // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(N2:: :vv)                           // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(N1:: :vv)                           // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(N1::aa) :vv)                        // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(ab):vv)                             // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(mapper(ab):arr[0:2])                       // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(mapper(aa) a:vv)                           // expected-warning {{missing ':' after ) - ignoring}}
-#pragma omp target update to(mapper(aa):vv)
+#pragma omp target update to(mapper(aa):vv) to(mapper(aa):arr[0])
 #pragma omp target update to(mapper(N1::stack<int>::id) :vv)
 
 #pragma omp target update from(mapper)                                  // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper()                                 // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper:vv)                               // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(:vv)                              // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(aa :vv)                           // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(N2:: :vv)                         // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(N1:: :vv)                         // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(N1::aa) :vv)                      // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(ab):vv)                           // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(mapper(ab):arr[0:2])                     // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(mapper(aa) a:vv)                         // expected-warning {{missing ':' after ) - ignoring}}
-#pragma omp target update from(mapper(aa):vv)
+#pragma omp target update from(mapper(aa):vv) from(mapper(aa):arr[0])
 #pragma omp target update from(mapper(N1::stack<int>::id) :vv)
     }
 #pragma omp declare mapper(id: vec v) map(v.len)                        // expected-error {{redefinition of user-defined mapper for type 'vec' with name 'id'}}
   }
   return arg;
 }