[OPENMP] Add support for explicit mapping of classes using 'this' pointer

[patricklyster]

Summary

Add support for explicit mapping of this pointer in OpenMP 5.0. Example use case:

class S {
  int a;
  void foo() {
    #pragma omp target map (this[0])
      a = 1;
  }
};

[ABataev]

LG

clang/include/clang/Basic/DiagnosticSemaKinds.td

def warn_omp_non_trivial_type_mapped : Warning<
   "Non-trivial type %0 is mapped, only trivial types are guaranteed to be mapped correctly">,
   InGroup<OpenMPTarget>;
 def err_omp_requires_clause_redeclaration : Error <
   "Only one %0 clause can appear on a requires directive in a single translation unit">;
 def note_omp_requires_previous_clause : Note <
   "%0 clause previously used here">;
 def err_omp_invalid_scope : Error <
   "'#pragma omp %0' directive must appear only in file scope">;
+def note_omp_invalid_length_on_this_ptr_mapping : Note <
+  "expected length on mapping of 'this' array section expression to be '1'">;
+def note_omp_invalid_lower_bound_on_this_ptr_mapping : Note <
+  "expected lower bound on mapping of 'this' array section expression to be '0' or not specified">;
+def note_omp_invalid_subscript_on_this_ptr_map : Note <
+  "expected 'this' subscript expression on map clause to be 'this[0]'">;
+def err_omp_invalid_map_this_expr : Error <
+  "invalid 'this' expression on 'map' clause">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForComponentList(
     auto CI = Components.rbegin();
     auto CE = Components.rend();
     auto I = CI;
 
     // Track if the map information being generated is the first for a list of
     // components.
     bool IsExpressionFirstInfo = true;
     Address BP = Address::invalid();
+    const Expr *AssocExpr = I->getAssociatedExpression();
+    const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
+    const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);   
 
-    if (isa<MemberExpr>(I->getAssociatedExpression())) {
+    if (isa<MemberExpr>(AssocExpr)) {
       // The base is the 'this' pointer. The content of the pointer is going
       // to be the base of the field being mapped.
       BP = CGF.LoadCXXThisAddress();
+    } else if ((AE && isa<CXXThisExpr>(AE->getBase()->IgnoreParenImpCasts())) ||
+               (OASE &&
+                isa<CXXThisExpr>(OASE->getBase()->IgnoreParenImpCasts()))) {
+      BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress();
     } else {
       // The base is the reference to the variable.
       // BP = &Var.
-      BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getAddress();
+      BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress();
       if (const auto *VD =
               dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) {
         if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
                 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
           if (*Res == OMPDeclareTargetDeclAttr::MT_Link) {
             IsLink = true;
             BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
           }

clang/lib/Sema/SemaOpenMP.cpp

 #include "clang/AST/ASTMutationListener.h"
 #include "clang/AST/CXXInheritance.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/StmtVisitor.h"
+#include "clang/AST/TypeOrdering.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
 #include "llvm/ADT/PointerEmbeddedInt.h"
 using namespace clang;
struct SharingMapTy {
     unsigned AssociatedLoops = 1;
     const Decl *PossiblyLoopCounter = nullptr;
     bool NowaitRegion = false;
     bool CancelRegion = false;
     bool LoopStart = false;
     SourceLocation InnerTeamsRegionLoc;
     /// Reference to the taskgroup task_reduction reference expression.
     Expr *TaskgroupReductionRef = nullptr;
+    llvm::DenseSet<QualType> MappedClassesQualTypes;
     SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name,
                  Scope *CurScope, SourceLocation Loc)
         : Directive(DKind), DirectiveName(Name), CurScope(CurScope),
           ConstructLoc(Loc) {}
     SharingMapTy() = default;
   };
 
   using StackTy = SmallVector<SharingMapTy, 4>;
getDoacrossDependClauses() const {
     const SharingMapTy &StackElem = Stack.back().first.back();
     if (isOpenMPWorksharingDirective(StackElem.Directive)) {
       const DoacrossDependMapTy &Ref = StackElem.DoacrossDepends;
       return llvm::make_range(Ref.begin(), Ref.end());
     }
     return llvm::make_range(StackElem.DoacrossDepends.end(),
                             StackElem.DoacrossDepends.end());
   }
+
+  // Store types of classes which have been explicitly mapped
+  void addMappedClassesQualTypes(QualType QT) {
+    SharingMapTy &StackElem = Stack.back().first.back();
+    StackElem.MappedClassesQualTypes.insert(QT);
+  }
+
+  // Return set of mapped classes types
+  bool isClassPreviouslyMapped(QualType QT) const {
+    const SharingMapTy &StackElem = Stack.back().first.back();
+    return StackElem.MappedClassesQualTypes.count(QT) != 0;
+  }
+
 };
 bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) {
   return isOpenMPParallelDirective(DKind) || isOpenMPTaskingDirective(DKind) ||
          isOpenMPTeamsDirective(DKind) || DKind == OMPD_unknown;
 }
 
 } // namespace
 
void VisitDeclRefExpr(DeclRefExpr *E) {
     }
   }
   void VisitMemberExpr(MemberExpr *E) {
     if (E->isTypeDependent() || E->isValueDependent() ||
         E->containsUnexpandedParameterPack() || E->isInstantiationDependent())
       return;
     auto *FD = dyn_cast<FieldDecl>(E->getMemberDecl());
     OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
-    if (isa<CXXThisExpr>(E->getBase()->IgnoreParens())) {
+    if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParens())) {
       if (!FD)
         return;
       DSAStackTy::DSAVarData DVar = Stack->getTopDSA(FD, /*FromParent=*/false);
       // Check if the variable has explicit DSA set and stop analysis if it
       // so.
       if (DVar.RefExpr || !ImplicitDeclarations.insert(FD).second)
         return;
 
if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParens())) {
                         ->getBase()
                         ->IgnoreParens());
               })) {
         // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
         //  A bit-field cannot appear in a map clause.
         //
         if (FD->isBitField())
           return;
+
+        // Check to see if the member expression is referencing a class that
+        // has already been explicitly mapped
+        if (Stack->isClassPreviouslyMapped(TE->getType()))
+          return;
+        
         ImplicitMap.emplace_back(E);
         return;
       }
 
       SourceLocation ELoc = E->getExprLoc();
       // OpenMP [2.9.3.6, Restrictions, p.2]
       //  A list item that appears in a reduction clause of the innermost
       //  enclosing worksharing or parallel construct may not be accessed in
if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
 
       // If we got an array subscript that express the whole dimension we
       // can have any array expressions before. If it only expressing part of
       // the dimension, we can only have unitary-size array expressions.
       if (checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE,
                                                       E->getType()))
         AllowWholeSizeArraySection = false;
 
+      if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
+        Expr::EvalResult Result;
+        if (CurE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) {
+          if (!Result.Val.getInt().isNullValue()) {
+            SemaRef.Diag(CurE->getIdx()->getExprLoc(),
+                         diag::err_omp_invalid_map_this_expr);
+            SemaRef.Diag(CurE->getIdx()->getExprLoc(),
+                         diag::note_omp_invalid_subscript_on_this_ptr_map);
+          }
+        }
+        RelevantExpr = TE;
+      }
+
       // Record the component - we don't have any declaration associated.
       CurComponents.emplace_back(CurE, nullptr);
     } else if (auto *CurE = dyn_cast<OMPArraySectionExpr>(E)) {
       assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
       E = CurE->getBase()->IgnoreParenImpCasts();
 
       QualType CurType =
           OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
         // 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)
             << CurE->getSourceRange();
         return nullptr;
       }
 
+      if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
+        Expr::EvalResult ResultR;
+        Expr::EvalResult ResultL;
+        if (CurE->getLength()->EvaluateAsInt(ResultR,
+                                             SemaRef.getASTContext())) {
+          if (!ResultR.Val.getInt().isOneValue()) {
+            SemaRef.Diag(CurE->getLength()->getExprLoc(),
+                         diag::err_omp_invalid_map_this_expr);
+            SemaRef.Diag(CurE->getLength()->getExprLoc(),
+                         diag::note_omp_invalid_length_on_this_ptr_mapping);
+          }
+        }
+        if (CurE->getLowerBound() && CurE->getLowerBound()->EvaluateAsInt(
+                                        ResultL, SemaRef.getASTContext())) {
+          if (!ResultL.Val.getInt().isNullValue()) {
+            SemaRef.Diag(CurE->getLowerBound()->getExprLoc(),
+                         diag::err_omp_invalid_map_this_expr);
+            SemaRef.Diag(CurE->getLowerBound()->getExprLoc(),
+                         diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
+          }
+        }
+        RelevantExpr = TE;
+      }
+
       // Record the component - we don't have any declaration associated.
       CurComponents.emplace_back(CurE, nullptr);
     } else {
       if (!NoDiagnose) {
         // If nothing else worked, this is not a valid map clause expression.
         SemaRef.Diag(
             ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
             << ERange;
for (Expr *RE : MVLI.VarList) {
     const Expr *BE = checkMapClauseExpressionBase(
         SemaRef, SimpleExpr, CurComponents, CKind, /*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());
+      // 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);
+      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.");
     assert(
         CurDeclaration->isCanonicalDecl() &&
         "Expecting components to have associated only canonical declarations.");

clang/test/OpenMP/target_ast_print.cpp

 // CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 
+// CHECK-LABEL: class S {
+class S {
+  void foo() {
+// CHECK-NEXT: void foo() {
+    int a = 0;
+// CHECK-NEXT: int a = 0;
+    #pragma omp target map(this[0])
+// CHECK-NEXT: #pragma omp target map(tofrom: this[0])
+      a++;
+// CHECK-NEXT: a++;
+    #pragma omp target map(this[:1])
+// CHECK-NEXT: #pragma omp target map(tofrom: this[:1])
+      a++;
+// CHECK-NEXT: a++;
+    #pragma omp target map((this)[0])
+// CHECK-NEXT: #pragma omp target map(tofrom: (this)[0])
+      a++;
+// CHECK-NEXT: a++;
+    #pragma omp target map(this[:a])
+// CHECK-NEXT: #pragma omp target map(tofrom: this[:a])
+      a++;
+// CHECK-NEXT: a++;
+    #pragma omp target map(this[a:1])
+// CHECK-NEXT: #pragma omp target map(tofrom: this[a:1])
+      a++;
+// CHECK-NEXT: a++;
+    #pragma omp target map(this[a])
+// CHECK-NEXT: #pragma omp target map(tofrom: this[a])
+      a++;
+// CHECK-NEXT: a++;
+  }
+// CHECK-NEXT: }
+};
+// CHECK-NEXT: };
+
 // CHECK-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
   int i, j, a[20], always;
 // CHECK-NEXT: int i, j, a[20]
 #pragma omp target
 // CHECK-NEXT: #pragma omp target
   foo();
 // CHECK-NEXT: foo();

clang/test/OpenMP/target_codegen.cpp

 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
 
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
 // CHECK-DAG: [[S1:%.+]] = type { double }
+// CHECK-DAG: [[S2:%.+]] = type { i32, i32, i32 }
 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
 // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
 // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
 
 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
 
 // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat
 
-// We have 8 target regions, but only 7 that actually will generate offloading
-// code and have mapped arguments, and only 5 have all-constant map sizes.
+// We have 9 target regions, but only 8 that actually will generate offloading
+// code and have mapped arguments, and only 6 have all-constant map sizes.
 
 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4]
 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
 // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ]] 2]
 // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i64] [i64 800]
 // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2]
 // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
 // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 288, i64 547, i64 547, i64 288, i64 288, i64 547, i64 547]
 // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]
 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 288, i64 288, i64 547]
+// CHECK-DAG: [[SIZET9:@.+]] = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 12]
+// CHECK-DAG: [[MAPT10:@.+]] = private unnamed_addr constant [1 x i64] [i64 35]
+// CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
 
 // Check if offloading descriptor is created.
 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
 // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8
 // CHECK: [[DEVEND:@.+]] = extern_weak constant i8
 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])

 // CHECK:       [[END]]
 // CHECK-NEXT:  br label %[[IFEND:.+]]
 // CHECK:       [[IFELSE]]
 // CHECK:       call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 // CHECK-NEXT:  br label %[[IFEND]]
 
 // CHECK:       [[IFEND]]
 
+// CHECK: define {{.*}}@{{.*}}zee{{.*}}
+
+// CHECK:       [[LOCAL_THIS:%.+]] = alloca [[S2]]*
+// CHECK:       [[BP:%.+]] = alloca [1 x i8*]
+// CHECK:       [[P:%.+]] = alloca [1 x i8*]
+// CHECK:       [[LOCAL_THIS1:%.+]] = load [[S2]]*, [[S2]]** [[LOCAL_THIS]]
+// CHECK:       [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], [[S2]]* [[LOCAL_THIS1]], i[[SZ]] 0
+// CHECK:       [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], [[S2]]* [[LOCAL_THIS1]], i[[SZ]] 0
+
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0    
+// CHECK-DAG:   [[PADDR0:%.+]] =  getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
+// CHECK-DAG:   [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to [[S2]]**
+// CHECK-DAG:   [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to [[S2]]**
+// CHECK-DAG:   store [[S2]]* [[ARR_IDX]], [[S2]]** [[CBPADDR0]]
+// CHECK-DAG:   store [[S2]]* [[ARR_IDX2]], [[S2]]** [[CPADDR0]]
+
+// CHECK:       [[BPR:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
+// CHECK:       [[PR:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
+// CHECK:       [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET9]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPT10]], i32 0, i32 0))
+// CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+// CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT0:@.+]]([[S2]]* [[LOCAL_THIS1]])
+// CHECK-NEXT:  br label %[[END]]
+// CHECK:       [[END]]
 
 // Check that the offloading functions are emitted and that the arguments are
 // correct and loaded correctly for the target regions of the callees of bar().
 
 // CHECK:       define internal void [[HVT7]]
 // Create local storage for each capture.
 // CHECK:       [[LOCAL_THIS:%.+]] = alloca [[S1]]*
 // CHECK:       [[LOCAL_B:%.+]] = alloca i[[SZ]]

 // CHECK-DAG:   load i16, i16* [[REF_AA]]
 // CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
 
 void bar () {
 #define pragma_target _Pragma("omp target")
 pragma_target
 {}
 }
+
+class S2 {
+  int a, b, c;
+
+public:
+  void zee() {
+    #pragma omp target map(this[0])
+      a++;
+  }
+};
+
+int main () {
+  S2 bar;
+  bar.zee();
+}
+
 #endif

clang/test/OpenMP/target_messages.cpp

 #pragma omp target
   ;
 #endif
 }
 #else
 void foo() {
 }
 
+class S {
+  public:
+  void zee() {
+    #pragma omp target map(this[:2]) // expected-note {{expected length on mapping of 'this' array section expression to be '1'}} // expected-error {{invalid 'this' expression on 'map' clause}}
+      int a;
+    #pragma omp target map(this[1:1]) // expected-note {{expected lower bound on mapping of 'this' array section expression to be '0' or null}} // expected-error {{invalid 'this' expression on 'map' clause}}
+      int b;
+    #pragma omp target map(this[1]) // expected-note {{expected 'this' subscript expression on map clause to be 'this[0]'}} // expected-error {{invalid 'this' expression on 'map' clause}}
+      int c;
+  }
+};
+
 #pragma omp target // expected-error {{unexpected OpenMP directive '#pragma omp target'}}
 
 int main(int argc, char **argv) {
   #pragma omp target { // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
   foo();
   #pragma omp target ( // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
   foo();
   #pragma omp target [ // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}