[OpenMP][Clang] Use ATTACH map-type for list-items with base-pointers.

clang/docs/ReleaseNotes.rst

@@ -619,6 +619,18 @@ OpenMP Support
 
 Improvements
 ^^^^^^^^^^^^
+- Mapping of expressions that have base-pointers now conforms to the OpenMP's
+  conditional pointer-attachment based on both pointee and poitner being
+  present, and one being new. This also lays the foundation of supporting
+  OpenMP 6.1's attach map-type modifier.
+- Several improvements were made to the handling of maps on list items involving
+  multiple levels of pointer dereferences, including not mapping intermediate
+  expressions, and grouping the items that share the same base-pointer, as
+  belonging to the same containing structure.
+- Support of array-sections on ``use_device_addr`` was made more robust,
+  including diagnosing when the array-section's base is not a named-variable.
+- Handling of ``use_device_addr`` and ``use_device_ptr`` in the presence of
+  other maps with the same base-pointer/variable, was improved.
 
 Additional Information
 ======================

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11702,6 +11702,9 @@ def err_omp_expected_addressable_lvalue_or_array_item : Error<
   "expected addressable lvalue expression, array element%select{ or array section|, array section or array shaping expression}0%select{| of non 'omp_depend_t' type}1">;
 def err_omp_expected_named_var_member_or_array_expression: Error<
   "expected expression containing only member accesses and/or array sections based on named variables">;
+def err_omp_expected_base_pointer_var_name_member_expr
+    : Error<"base-pointer is not a variable name%select{| or data member of "
+            "current class}0">;
 def err_omp_bit_fields_forbidden_in_clause : Error<
   "bit fields cannot be used to specify storage in a '%0' clause">;
 def err_array_section_does_not_specify_contiguous_storage : Error<

clang/lib/Sema/SemaOpenMP.cpp

@@ -2217,6 +2217,7 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
     // | ptr  |      n.a.     |  -  |   x   |       -       |     -    | bycopy|
     // | ptr  |      n.a.     |  x  |   -   |       -       |     -    | null  |
     // | ptr  |      n.a.     |  -  |   -   |       -       |     x    | byref |
+    // | ptr  |      n.a.     |  -  |   -   |       -       |   x, x[] | bycopy|
     // | ptr  |      n.a.     |  -  |   -   |       -       |    x[]   | bycopy|
     // | ptr  |      n.a.     |  -  |   -   |       x       |          | bycopy|
     // | ptr  |      n.a.     |  -  |   -   |       x       |     x    | bycopy|
@@ -2242,18 +2243,22 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
     //  - For pointers mapped by value that have either an implicit map or an
     //    array section, the runtime library may pass the NULL value to the
     //    device instead of the value passed to it by the compiler.
+    //  - If both a pointer an a dereference of it are mapped, then the pointer
+    //    should be passed by reference.
 
     if (Ty->isReferenceType())
       Ty = Ty->castAs<ReferenceType>()->getPointeeType();
 
-    // Locate map clauses and see if the variable being captured is referred to
-    // in any of those clauses. Here we only care about variables, not fields,
-    // because fields are part of aggregates.
+    // Locate map clauses and see if the variable being captured is mapped by
+    // itself, or referred to, in any of those clauses. Here we only care about
+    // variables, not fields, because fields are part of aggregates.
     bool IsVariableAssociatedWithSection = false;
+    bool IsVariableItselfMapped = false;
 
     DSAStack->checkMappableExprComponentListsForDeclAtLevel(
         D, Level,
         [&IsVariableUsedInMapClause, &IsVariableAssociatedWithSection,
+         &IsVariableItselfMapped,
          D](OMPClauseMappableExprCommon::MappableExprComponentListRef
                 MapExprComponents,
             OpenMPClauseKind WhereFoundClauseKind) {
@@ -2269,8 +2274,19 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
 
           assert(EI != EE && "Invalid map expression!");
 
-          if (isa<DeclRefExpr>(EI->getAssociatedExpression()))
-            IsVariableUsedInMapClause |= EI->getAssociatedDeclaration() == D;
+          if (isa<DeclRefExpr>(EI->getAssociatedExpression()) &&
+              EI->getAssociatedDeclaration() == D) {
+            IsVariableUsedInMapClause = true;
+
+            // If the component list has only one element, it's for mapping the
+            // variable itself, like map(p). This takes precedence in
+            // determining how it's captured, so we don't need to look further
+            // for any other maps that use the variable (like map(p[0]) etc.)
+            if (MapExprComponents.size() == 1) {
+              IsVariableItselfMapped = true;
+              return true;
+            }
+          }
 
           ++EI;
           if (EI == EE)
@@ -2284,8 +2300,10 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
               isa<MemberExpr>(EI->getAssociatedExpression()) ||
               isa<OMPArrayShapingExpr>(Last->getAssociatedExpression())) {
             IsVariableAssociatedWithSection = true;
-            // There is nothing more we need to know about this variable.
-            return true;
+            // We've found a case like map(p[0]) or map(p->a) or map(*p),
+            // so we are done with this particular map, but we need to keep
+            // looking in case we find a map(p).
+            return false;
           }
 
           // Keep looking for more map info.
@@ -2294,8 +2312,23 @@ bool SemaOpenMP::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
 
     if (IsVariableUsedInMapClause) {
       // If variable is identified in a map clause it is always captured by
-      // reference except if it is a pointer that is dereferenced somehow.
-      IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection);
+      // reference except if it is a pointer that is dereferenced somehow, but
+      // not itself mapped.
+      //
+      // OpenMP 6.0, 7.1.1: Data sharing attribute rules, variables referenced
+      // in a construct::
+      // If a list item in a has_device_addr clause or in a map clause on the
+      // target construct has a base pointer, and the base pointer is a scalar
+      // variable *that is not a list item in a map clause on the construct*,
+      // the base pointer is firstprivate.
+      //
+      // OpenMP 4.5, 2.15.1.1: Data-sharing Attribute Rules for Variables
+      // Referenced in a Construct:
+      // If an array section is a list item in a map clause on the target
+      // construct and the array section is derived from a variable for which
+      // the type is pointer then that variable is firstprivate.
+      IsByRef = IsVariableItselfMapped ||
+                !(Ty->isPointerType() && IsVariableAssociatedWithSection);
     } else {
       // By default, all the data that has a scalar type is mapped by copy
       // (except for reduction variables).
@@ -22838,8 +22871,10 @@ static void checkMappableExpressionList(
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown,
     ArrayRef<OpenMPMapModifierKind> Modifiers = {},
     bool IsMapTypeImplicit = false, bool NoDiagnose = false) {
-  // We only expect mappable expressions in 'to', 'from', and 'map' clauses.
-  assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) &&
+  // We only expect mappable expressions in 'to', 'from', 'map', and
+  // 'use_device_addr' clauses.
+  assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from ||
+          CKind == OMPC_use_device_addr) &&
          "Unexpected clause kind with mappable expressions!");
   unsigned OMPVersion = SemaRef.getLangOpts().OpenMP;
 
@@ -24543,17 +24578,67 @@ SemaOpenMP::ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
     // similar properties of a first private variable.
     DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
 
-    // Create a mappable component for the list item. List items in this clause
-    // only need a component.
-    MVLI.VarBaseDeclarations.push_back(D);
-    MVLI.VarComponents.emplace_back();
-    Expr *Component = SimpleRefExpr;
-    if (VD && (isa<ArraySectionExpr>(RefExpr->IgnoreParenImpCasts()) ||
-               isa<ArraySubscriptExpr>(RefExpr->IgnoreParenImpCasts())))
-      Component =
-          SemaRef.DefaultFunctionArrayLvalueConversion(SimpleRefExpr).get();
-    MVLI.VarComponents.back().emplace_back(Component, D,
-                                           /*IsNonContiguous=*/false);
+    // Use the map-like approach to fully populate VarComponents
+    OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
+
+    const Expr *BE = checkMapClauseExpressionBase(
+        SemaRef, RefExpr, CurComponents, OMPC_use_device_addr,
+        DSAStack->getCurrentDirective(),
+        /*NoDiagnose=*/false);
+
+    if (!BE)
+      continue;
+
+    assert(!CurComponents.empty() &&
+           "use_device_addr clause expression with no components!");
+
+    // OpenMP use_device_addr: If a list item is an array section, the array
+    // base must be a base language identifier. We caught the cases where
+    // the array-section has a base-variable in getPrivateItem. e.g.
+    //   struct S {
+    //     int a[10];
+    //   }; S s1;
+    //   ... use_device_addr(s1.a[0])    // not ok, caught already
+    //
+    // But we still neeed to verify that the base-pointer is also a
+    // base-language identifier, and catch cases like:
+    //   int *pa[10]; *p;
+    //   ... use_device_addr(pa[1][2])   // not ok, base-pointer is pa[1]
+    //   ... use_device_addr(p[1])       // ok
+    //   ... use_device_addr(this->p[1]) // ok
+    auto AttachPtrResult = OMPClauseMappableExprCommon::findAttachPtrExpr(
+        CurComponents, DSAStack->getCurrentDirective());
+    const Expr *AttachPtrExpr = AttachPtrResult.first;
+
+    if (AttachPtrExpr) {
+      const Expr *BaseExpr = AttachPtrExpr->IgnoreParenImpCasts();
+      bool IsValidBase = false;
+
+      if (isa<DeclRefExpr>(BaseExpr))
+        IsValidBase = true;
+      else if (const auto *ME = dyn_cast<MemberExpr>(BaseExpr);
+               ME && isa<CXXThisExpr>(ME->getBase()->IgnoreParenImpCasts()))
+        IsValidBase = true;
+
+      if (!IsValidBase) {
+        SemaRef.Diag(ELoc,
+                     diag::err_omp_expected_base_pointer_var_name_member_expr)
+            << (SemaRef.getCurrentThisType().isNull() ? 0 : 1)
+            << AttachPtrExpr->getSourceRange();
+        continue;
+      }
+    }
+
+    // Get the declaration from the components
+    ValueDecl *CurDeclaration = CurComponents.back().getAssociatedDeclaration();
+    assert(isa<CXXThisExpr>(BE) ||
+           CurDeclaration &&
+               "Unexpected null decl for use_device_addr clause.");
+
+    MVLI.VarBaseDeclarations.push_back(CurDeclaration);
+    MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+    MVLI.VarComponents.back().append(CurComponents.begin(),
+                                     CurComponents.end());
   }
 
   if (MVLI.ProcessedVarList.empty())

clang/test/OpenMP/bug59160.c

@@ -8,7 +8,13 @@ void zoo(void) {
   xp[1] = &x[0];
   short **xpp = &xp[0];
   x[1] = 111;
-#pragma omp target data map(tofrom: xpp[1][1]) use_device_addr(xpp[1][1])
+
+// NOTE: use_device_addr on xpp[1][1] is non-compliant, as the base-pointer
+// is xpp[1], which is not a base-language identifier.
+#pragma omp target data map(tofrom: xpp[1][1]) //use_device_addr(xpp[1][1])
+// FIXME: The assumption that xpp should not be mapped is incorrect.
+// The base-pointer of the array-section is xpp[1], not xpp, so the implicit
+// clause on xpp, i.e. a zero-length array-section amp, should still be emitted.
 #pragma omp target has_device_addr(xpp[1][1])
   {
     xpp[1][1] = 222;