Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7703,15 +7703,18 @@ dyn_cast(I->getAssociatedExpression()); const auto *OAShE = dyn_cast(I->getAssociatedExpression()); - const auto *UO = dyn_cast(I->getAssociatedExpression()); - const auto *BO = dyn_cast(I->getAssociatedExpression()); bool IsPointer = OAShE || (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE) .getCanonicalType() ->isAnyPointerType()) || I->getAssociatedExpression()->getType()->isAnyPointerType(); - bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous; + bool IsNonDerefPointer = + IsPointer && !isa(I->getAssociatedExpression()) && + !isa(I->getAssociatedExpression()) && + !isa(I->getAssociatedExpression()) && + !isa(I->getAssociatedExpression()) && + !IsNonContiguous; if (OASE) ++DimSize; @@ -7725,7 +7728,9 @@ isa(Next->getAssociatedExpression()) || isa(Next->getAssociatedExpression()) || isa(Next->getAssociatedExpression()) || - isa(Next->getAssociatedExpression())) && + isa(Next->getAssociatedExpression()) || + isa( + Next->getAssociatedExpression())) && "Unexpected expression"); Address LB = Address::invalid(); Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -16837,7 +16837,12 @@ public: bool VisitDeclRefExpr(DeclRefExpr *DRE) { - if (!isa(DRE->getDecl())) { + // Allow variables, and also allow functions if they are part of a call + // expression. However, we must not map functions themselves. This is + // prevented since a call expression will always be added to Components + // first. + if (!isa(DRE->getDecl()) && + (Components.empty() || !isa(DRE->getDecl()))) { emitErrorMsg(); return false; } @@ -16860,6 +16865,13 @@ E = BaseE; } + // Allow results of method calls to be mapped. + if (isa(ME->getMemberDecl())) { + RelevantExpr = ME; + Components.emplace_back(ME, ME->getMemberDecl(), IsNonContiguous); + return true; + } + if (!isa(ME->getMemberDecl())) { if (!NoDiagnose) { SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field) @@ -16920,6 +16932,13 @@ return RelevantExpr || Visit(E); } + bool VisitOpaqueValueExpr(OpaqueValueExpr *OVE) { + // An OpaqueValue expression may be encountered e.g. in the operands of a + // binary conditional ("Elvis") operator. For now, it is sufficient to + // forward to the source expression. + return Visit(OVE->getSourceExpr()->IgnoreParenImpCasts()); + } + bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) { Expr *E = AE->getBase()->IgnoreParenImpCasts(); @@ -17050,8 +17069,9 @@ } bool VisitUnaryOperator(UnaryOperator *UO) { - if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() || - UO->getOpcode() != UO_Deref) { + if (SemaRef.getLangOpts().OpenMP < 50 || + (Components.empty() && + (!UO->isLValue() || UO->getOpcode() != UO_Deref))) { emitErrorMsg(); return false; } @@ -17092,6 +17112,36 @@ Components.emplace_back(COCE, nullptr, IsNonContiguous); return true; } + bool VisitCallExpr(CallExpr *CE) { + if (SemaRef.getLangOpts().OpenMP < 50 || + (Components.empty() && !CE->isLValue())) { + emitErrorMsg(); + return false; + } + assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); + Components.emplace_back(CE, nullptr, IsNonContiguous); + return RelevantExpr || Visit(CE->getCallee()->IgnoreParenImpCasts()); + } + bool VisitCastExpr(CastExpr *CE) { + if (SemaRef.getLangOpts().OpenMP < 50 || + (Components.empty() && !CE->isLValue())) { + emitErrorMsg(); + return false; + } + assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); + Components.emplace_back(CE, nullptr, false); + return RelevantExpr || Visit(CE->getSubExpr()->IgnoreParenImpCasts()); + } + bool VisitAbstractConditionalOperator(AbstractConditionalOperator *ACO) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); + Components.emplace_back(ACO, nullptr, false); + return RelevantExpr || Visit(ACO->getTrueExpr()->IgnoreParenImpCasts()) || + Visit(ACO->getFalseExpr()->IgnoreParenImpCasts()); + } bool VisitStmt(Stmt *) { emitErrorMsg(); return false; @@ -17268,7 +17318,7 @@ } QualType DerivedType = - std::prev(CI)->getAssociatedDeclaration()->getType(); + std::prev(CI)->getAssociatedExpression()->getType(); SourceLocation DerivedLoc = std::prev(CI)->getAssociatedExpression()->getExprLoc(); @@ -17277,6 +17327,20 @@ // will be considered to be T for all purposes of this clause. DerivedType = DerivedType.getNonReferenceType(); + // There are restrictions on when arrays and structs may be mapped which + // are checked in this function. It is difficult to check these + // conditions at compile time for a general lvalue (e.g. for a ternary), + // so instead we whitelist the kind of expressions we would like to + // perform these checks on. + const Expr *nextComponentExpr = (CI == CE) + ? SI->getAssociatedExpression() + : CI->getAssociatedExpression(); + const bool shouldCheckForOverlap = + isa(nextComponentExpr) || + isa(nextComponentExpr) || + isa(nextComponentExpr) || + isa(nextComponentExpr); + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.1] // A variable for which the type is pointer and an array section // derived from that variable must not appear as list items of map @@ -17288,7 +17352,7 @@ // storage in the device data environment, all of the original storage // must have corresponding storage in the device data environment. // - if (DerivedType->isAnyPointerType()) { + if (shouldCheckForOverlap && DerivedType->isAnyPointerType()) { if (CI == CE || SI == SE) { SemaRef.Diag( DerivedLoc, @@ -17316,7 +17380,8 @@ // original storage. // // An expression is a subset of the other. - if (CurrentRegionOnly && (CI == CE || SI == SE)) { + if (shouldCheckForOverlap && CurrentRegionOnly && + (CI == CE || SI == SE)) { if (CKind == OMPC_map) { if (CI != CE || SI != SE) { // Allow constructs like this: map(s, s.ptr[0:1]), where s.ptr is @@ -17351,7 +17416,7 @@ // The current expression uses the same base as other expression in the // data environment but does not contain it completely. - if (!CurrentRegionOnly && SI != SE) + if (shouldCheckForOverlap && !CurrentRegionOnly && SI != SE) EnclosingExpr = RE; // The current expression is a subset of the expression in the data @@ -17653,8 +17718,10 @@ auto *VD = dyn_cast(CurDeclaration); const auto *FD = dyn_cast(CurDeclaration); + const auto *FnD = dyn_cast(CurDeclaration); - assert((VD || FD) && "Only variables or fields are expected here!"); + assert((VD || FD || FnD) && + "Only variables, fields, or functions are expected here!"); (void)FD; // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10] Index: clang/test/OpenMP/target_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_map_messages.cpp +++ clang/test/OpenMP/target_map_messages.cpp @@ -1,16 +1,16 @@ -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized - -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 400 -x c %s -Wno-openmp -Wuninitialized + +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=40 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 400 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 400 -x c %s -Wno-openmp-mapping -Wuninitialized #ifdef CCODE void foo(int arg) { const int n = 0; @@ -42,6 +42,14 @@ SREF(int &a) : a(a) {} }; +int &id(int &x) { + return x; +} + +int *id2(int *x) { + return x; +} + template struct SA { static int ss; @@ -73,117 +81,163 @@ {} #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}} {} +#pragma omp target map(id) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(id(b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(id(id(b[3]))) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(id2(&b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(*id2(&b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} - #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}} +#pragma omp target map(to \ + : ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}} {} - #pragma omp target map(to:b,e) +#pragma omp target map(to \ + : b, e) {} - #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(to \ + : b, e) map(to \ + : b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} {} - #pragma omp target map(to:b[:2],e) +#pragma omp target map(to \ + : b[:2], e) {} - #pragma omp target map(to:b,e[:]) +#pragma omp target map(to \ + : b, e[:]) {} - #pragma omp target map(b[-1:]) // expected-error {{array section must be a subset of the original array}} +#pragma omp target map(b [-1:]) // expected-error {{array section must be a subset of the original array}} {} - #pragma omp target map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}} +#pragma omp target map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}} {} - #pragma omp target map(b[true:true]) +#pragma omp target map(b [true:true]) {} - #pragma omp target map(: c,f) // expected-error {{missing map type}} +#pragma omp target map( \ + : c, f) // expected-error {{missing map type}} {} - #pragma omp target map(always, tofrom: c,f) +#pragma omp target map(always, tofrom \ + : c, f) {} - #pragma omp target map(always, tofrom: c[1:2],f) +#pragma omp target map(always, tofrom \ + : c [1:2], f) {} - #pragma omp target map(always, tofrom: c,f[1:2]) +#pragma omp target map(always, tofrom \ + : c, f [1:2]) {} - #pragma omp target map(always, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(always, tofrom \ + : c[:], f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(always, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(always, tofrom \ + : c, f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(always) // expected-error {{use of undeclared identifier 'always'}} +#pragma omp target map(always) // expected-error {{use of undeclared identifier 'always'}} {} - #pragma omp target map(close, tofrom: c,f) +#pragma omp target map(close, tofrom \ + : c, f) {} - #pragma omp target map(close, tofrom: c[1:2],f) +#pragma omp target map(close, tofrom \ + : c [1:2], f) {} - #pragma omp target map(close, tofrom: c,f[1:2]) +#pragma omp target map(close, tofrom \ + : c, f [1:2]) {} - #pragma omp target map(close, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(close, tofrom \ + : c[:], f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(close, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target map(close, tofrom \ + : c, f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}} +#pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}} {} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, tofrom: c,f) +#pragma omp target map(present, tofrom \ + : c, f) {} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, tofrom: c[1:2],f) +#pragma omp target map(present, tofrom \ + : c [1:2], f) {} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, tofrom: c,f[1:2]) +#pragma omp target map(present, tofrom \ + : c, f [1:2]) {} // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, tofrom: c[:],f) +#pragma omp target map(present, tofrom \ + : c[:], f) {} // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, tofrom: c,f[:]) +#pragma omp target map(present, tofrom \ + : c, f[:]) {} // expected-error@+1 {{use of undeclared identifier 'present'}} - #pragma omp target map(present) +#pragma omp target map(present) {} - #pragma omp target map(close, close, tofrom: a) // expected-error {{same map type modifier has been specified more than once}} +#pragma omp target map(close, close, tofrom \ + : a) // expected-error {{same map type modifier has been specified more than once}} {} - #pragma omp target map(always, close, always, close, tofrom: a) // expected-error 2 {{same map type modifier has been specified more than once}} +#pragma omp target map(always, close, always, close, tofrom \ + : a) // expected-error 2 {{same map type modifier has been specified more than once}} {} // ge51-error@+2 {{same map type modifier has been specified more than once}} // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(present, present, tofrom: a) +#pragma omp target map(present, present, tofrom \ + : a) {} // expected-error@+3 2 {{same map type modifier has been specified more than once}} // ge51-error@+2 1 {{same map type modifier has been specified more than once}} // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(always, close, present, always, close, present, tofrom: a) +#pragma omp target map(always, close, present, always, close, present, tofrom \ + : a) {} - #pragma omp target map( , tofrom: a) // expected-error {{missing map type modifier}} +#pragma omp target map(, tofrom \ + : a) // expected-error {{missing map type modifier}} {} - #pragma omp target map( , , tofrom: a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} +#pragma omp target map(, , tofrom \ + : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} {} - #pragma omp target map( , , : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}} +#pragma omp target map(, , \ + : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}} {} // ge51-error@+3 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} // lt51-error@+2 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} // expected-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} - #pragma omp target map( d, f, bf: a) +#pragma omp target map(d, f, bf \ + : a) {} // expected-error@+4 {{missing map type modifier}} // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} // expected-error@+1 {{missing map type}} - #pragma omp target map( , f, : a) +#pragma omp target map(, f, \ + : a) {} - #pragma omp target map(always close: a) // expected-error {{missing map type}} +#pragma omp target map(always close \ + : a) // expected-error {{missing map type}} {} - #pragma omp target map(always close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(always close bf \ + : a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} {} // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} // expected-error@+1 {{missing map type}} - #pragma omp target map(always tofrom close: a) +#pragma omp target map(always tofrom close \ + : a) {} // ge51-error@+2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(tofrom from: a) +#pragma omp target map(tofrom from \ + : a) {} - #pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(close bf \ + : a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} {} - #pragma omp target map(([b[I]][bf])f) // lt50-error {{expected ',' or ']' in lambda capture list}} lt50-error {{expected ')'}} lt50-note {{to match this '('}} +#pragma omp target map(([b[I]][bf])f) // lt50-error {{expected ',' or ']' in lambda capture list}} lt50-error {{expected ')'}} lt50-note {{to match this '('}} {} return; } @@ -197,6 +251,9 @@ float *foo() { return &Arr[0]; } + float &bar() { + return Arr[0]; + } }; struct SC { @@ -356,7 +413,13 @@ {} #pragma omp target map(r.S.Arr[:12]) {} -#pragma omp target map(r.S.foo() [:12]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} +#pragma omp target map(r.S.foo) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(r.S.foo()) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(r.S.foo() [:12]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(r.S.bar()) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} #pragma omp target map(r.C, r.D) {} @@ -570,7 +633,10 @@ #pragma omp target map(to, x) foo(); #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -#pragma omp target data map(tofrom: argc > 0 ? x : y) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error 2 {{expected addressable lvalue in 'map' clause}} +#pragma omp target data map(tofrom \ + : argc > 0 ? x : y) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target data map(tofrom \ + : x ?: y) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target data map(argc) #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}} #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning 2 {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}} @@ -668,6 +734,11 @@ virtual void foo() {} } s9; +class S10 : public S9 { +public: + int a; + void foo() {} +} s10; int main(int argc, char **argv) { const int d = 5; @@ -704,7 +775,10 @@ #pragma omp target map(to, x) foo(); #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} +#pragma omp target data map(tofrom \ + : argc > 0 ? argv[1] : argv[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target data map(tofrom \ + : argv[1] ?: argv[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target data map(argc) #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}} #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}} @@ -835,6 +909,24 @@ {} #pragma omp target map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}} {} +#pragma omp target map(reinterpret_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(*reinterpret_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(static_cast (s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(*static_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(*dynamic_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map(const_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(*const_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target map((S9)s10) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target map(*((S9 *)&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} #pragma omp target map(delete: a) // expected-error {{map type 'delete' is not allowed for '#pragma omp target'}} {} Index: clang/test/OpenMP/target_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_map_messages.cpp +++ clang/test/OpenMP/target_parallel_for_map_messages.cpp @@ -127,11 +127,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel for map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel for map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(S1) // expected-error {{'S1' does not refer to a value}} @@ -242,11 +246,15 @@ #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel for map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel for map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp +++ clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp @@ -127,11 +127,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel for simd map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel for simd map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}} @@ -242,11 +246,15 @@ #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel for simd map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel for simd map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_parallel_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_map_messages.cpp +++ clang/test/OpenMP/target_parallel_map_messages.cpp @@ -127,11 +127,14 @@ foo(); #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel map(tofrom \ : argc > 0 ? x : y) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel map(tofrom \ + : x ?: y) + foo(); #pragma omp target parallel map(argc) foo(); #pragma omp target parallel map(S1) // expected-error {{'S1' does not refer to a value}} @@ -240,11 +243,14 @@ foo(); #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target parallel map(tofrom \ : argc > 0 ? argv[1] : argv[2]) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target parallel map(tofrom \ + : argv[1] ?: argv[2]) + foo(); #pragma omp target parallel map(argc) foo(); #pragma omp target parallel map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_simd_map_messages.cpp +++ clang/test/OpenMP/target_simd_map_messages.cpp @@ -121,11 +121,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target simd map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target simd map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(S1) // expected-error {{'S1' does not refer to a value}} @@ -230,11 +234,15 @@ #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target simd map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target simd map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_teams_distribute_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_map_messages.cpp @@ -127,11 +127,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(S1) // expected-error {{'S1' does not refer to a value}} @@ -242,11 +246,15 @@ #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp @@ -125,11 +125,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute parallel for map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute parallel for map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(S1) // expected-error {{'S1' does not refer to a value}} @@ -239,11 +243,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute parallel for map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute parallel for map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp @@ -127,11 +127,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute parallel for simd map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute parallel for simd map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}} @@ -241,11 +245,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute parallel for simd map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute parallel for simd map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp @@ -127,11 +127,15 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute simd map(tofrom \ : argc > 0 ? x : y) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute simd map(tofrom \ + : x ?: y) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(S1) // expected-error {{'S1' does not refer to a value}} @@ -242,11 +246,15 @@ #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -// ge50-error@+3 {{expected addressable lvalue in 'map' clause}} // lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target teams distribute simd map(tofrom \ : argc > 0 ? argv[1] : argv[2]) for (i = 0; i < argc; ++i) foo(); +// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams distribute simd map(tofrom \ + : argv[1] ?: argv[2]) + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute simd map(argc) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(S1) // expected-error {{'S1' does not refer to a value}} Index: clang/test/OpenMP/target_teams_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_map_messages.cpp +++ clang/test/OpenMP/target_teams_map_messages.cpp @@ -30,6 +30,14 @@ ; } +int &id(int &x) { + return x; +} + +int *id2(int *x) { + return x; +} + template struct SA { static int ss; @@ -56,34 +64,55 @@ {} #pragma omp target teams map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}} {} +#pragma omp target teams map(id) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(id(b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map(id(id(b[3]))) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map(id2(&b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*id2(&b[3])) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} - #pragma omp target teams map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}} +#pragma omp target teams map(to \ + : ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}} {} - #pragma omp target teams map(to:b,e) +#pragma omp target teams map(to \ + : b, e) {} - #pragma omp target teams map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams map(to \ + : b, e) map(to \ + : b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} {} - #pragma omp target teams map(to:b[:2],e) +#pragma omp target teams map(to \ + : b[:2], e) {} - #pragma omp target teams map(to:b,e[:]) +#pragma omp target teams map(to \ + : b, e[:]) {} - #pragma omp target teams map(b[-1:]) // expected-error {{array section must be a subset of the original array}} +#pragma omp target teams map(b [-1:]) // expected-error {{array section must be a subset of the original array}} {} - #pragma omp target teams map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}} +#pragma omp target teams map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}} {} - #pragma omp target teams map(b[true:true]) +#pragma omp target teams map(b [true:true]) {} - #pragma omp target teams map(always, tofrom: c,f) +#pragma omp target teams map(always, tofrom \ + : c, f) {} - #pragma omp target teams map(always, tofrom: c[1:2],f) +#pragma omp target teams map(always, tofrom \ + : c [1:2], f) {} - #pragma omp target teams map(always, tofrom: c,f[1:2]) +#pragma omp target teams map(always, tofrom \ + : c, f [1:2]) {} - #pragma omp target teams map(always, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target teams map(always, tofrom \ + : c[:], f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target teams map(always, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target teams map(always, tofrom \ + : c, f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} return; } @@ -97,6 +126,9 @@ float *foo() { return &Arr[0]; } + float &bar() { + return Arr[0]; + } }; struct SC { @@ -134,196 +166,209 @@ SD u; SC r(p),t(p); - #pragma omp target teams map(r) +#pragma omp target teams map(r) {} - #pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[2] [0:2] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:] [0:2] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[2][3][0:2]) +#pragma omp target teams map(marr[2][3] [0:2]) {} - #pragma omp target teams map(marr[:][:][:]) +#pragma omp target teams map(marr[:][:][:]) {} - #pragma omp target teams map(marr[:2][:][:]) +#pragma omp target teams map(marr[:2][:][:]) {} - #pragma omp target teams map(marr[arg:][:][:]) +#pragma omp target teams map(marr [arg:][:][:]) {} - #pragma omp target teams map(marr[arg:]) +#pragma omp target teams map(marr [arg:]) {} - #pragma omp target teams map(marr[arg:][:arg][:]) // correct if arg is the size of dimension 2 +#pragma omp target teams map(marr [arg:][:arg][:]) // correct if arg is the size of dimension 2 {} - #pragma omp target teams map(marr[:arg][:]) +#pragma omp target teams map(marr[:arg][:]) {} - #pragma omp target teams map(marr[:arg][n:]) +#pragma omp target teams map(marr[:arg] [n:]) {} - #pragma omp target teams map(marr[:][:arg][n:]) // correct if arg is the size of dimension 2 +#pragma omp target teams map(marr[:][:arg] [n:]) // correct if arg is the size of dimension 2 {} - #pragma omp target teams map(marr[:][:m][n:]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:][:m] [n:]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[n:m][:arg][n:]) +#pragma omp target teams map(marr [n:m][:arg] [n:]) {} - #pragma omp target teams map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:2][1:][:]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:2] [1:][:]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:2][:][1:]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:2][:] [1:]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:1][:2][:]) +#pragma omp target teams map(marr[:1][:2][:]) {} - #pragma omp target teams map(marr[:1][0][:]) +#pragma omp target teams map(marr[:1][0][:]) {} - #pragma omp target teams map(marr[:arg][:2][:]) // correct if arg is 1 +#pragma omp target teams map(marr[:arg][:2][:]) // correct if arg is 1 {} - #pragma omp target teams map(marr[:1][3:1][:2]) +#pragma omp target teams map(marr[:1] [3:1][:2]) {} - #pragma omp target teams map(marr[:1][3:arg][:2]) // correct if arg is 1 +#pragma omp target teams map(marr[:1] [3:arg][:2]) // correct if arg is 1 {} - #pragma omp target teams map(marr[:1][3:2][:2]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:1] [3:2][:2]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr[:2][:10][:]) +#pragma omp target teams map(marr[:2][:10][:]) {} - #pragma omp target teams map(marr[:2][:][:5+5]) +#pragma omp target teams map(marr[:2][:][:5 + 5]) {} - #pragma omp target teams map(marr[:2][2+2-4:][0:5+5]) +#pragma omp target teams map(marr[:2] [2 + 2 - 4:] [0:5 + 5]) {} - #pragma omp target teams map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(marr2[:1][:2][0]) +#pragma omp target teams map(marr2[:1][:2][0]) {} - #pragma omp target teams map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1. +#pragma omp target teams map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1. {} - #pragma omp target teams map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2. +#pragma omp target teams map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2. {} - #pragma omp target teams map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(mvla[1][2:arg][:]) +#pragma omp target teams map(mvla[1] [2:arg][:]) {} - #pragma omp target teams map(mvla[:1][:][:]) +#pragma omp target teams map(mvla[:1][:][:]) {} - #pragma omp target teams map(mvla2[:1][:2][:11]) +#pragma omp target teams map(mvla2[:1][:2][:11]) {} - #pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(mptr[:2] [2 + 2 - 4:1] [0:5 + 5]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(mptr[:1][:2-1][2:4-3]) +#pragma omp target teams map(mptr[:1][:2 - 1] [2:4 - 3]) {} - #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // correct if arg is 1. +#pragma omp target teams map(mptr[:1][:arg] [2:4 - 3]) // correct if arg is 1. {} - #pragma omp target teams map(mptr[:1][:2-1][0:2]) +#pragma omp target teams map(mptr[:1][:2 - 1] [0:2]) {} - #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(mptr[:1][:2] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target teams map(mptr[:1][:] [0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}} +#pragma omp target teams map(mptr[:2][:1] [0:2]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(r.ArrS[0].B) +#pragma omp target teams map(r.ArrS[0].B) {} - #pragma omp target teams map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target teams map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}} {} - #pragma omp target teams map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target teams map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}} {} - #pragma omp target teams map(r.ArrS[0].Arr[1:23]) +#pragma omp target teams map(r.ArrS[0].Arr [1:23]) {} - #pragma omp target teams map(r.ArrS[0].Arr[1:arg]) +#pragma omp target teams map(r.ArrS[0].Arr [1:arg]) {} - #pragma omp target teams map(r.ArrS[0].Arr[arg:23]) +#pragma omp target teams map(r.ArrS[0].Arr [arg:23]) {} - #pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}} +#pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}} {} - #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}} +#pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}} {} - #pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A) +#pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A) {} - #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} +#pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} {} - #pragma omp target teams map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target teams map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} {} - #pragma omp target teams map(r.PtrS->A, r.PtrS->B) +#pragma omp target teams map(r.PtrS->A, r.PtrS->B) {} - #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} +#pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} {} - #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} {} - #pragma omp target teams map(r.RPtrS->A, r.RPtrS->B) +#pragma omp target teams map(r.RPtrS->A, r.RPtrS->B) {} - #pragma omp target teams map(r.S.Arr[:12]) +#pragma omp target teams map(r.S.Arr[:12]) {} - // ge50-error@+2 {{expected addressable lvalue in 'map' clause}} - // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}} - #pragma omp target teams map(r.S.foo()[:12]) +#pragma omp target teams map(r.S.foo) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(r.S.foo()) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*(r.S.foo())) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams map(r.S.foo() [:12]) + {} +#pragma omp target teams map(r.S.bar()) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} - #pragma omp target teams map(r.C, r.D) +#pragma omp target teams map(r.C, r.D) {} - #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} {} - #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} {} - #pragma omp target teams map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else. +#pragma omp target teams map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else. {} - #pragma omp target teams map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} {} - #pragma omp target teams map(r.C, t.C) +#pragma omp target teams map(r.C, t.C) {} - #pragma omp target teams map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}} +#pragma omp target teams map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}} {} - #pragma omp target teams map(r.Arr) +#pragma omp target teams map(r.Arr) {} - #pragma omp target teams map(r.Arr[3:5]) +#pragma omp target teams map(r.Arr [3:5]) {} - #pragma omp target teams map(r.Ptr[3:5]) +#pragma omp target teams map(r.Ptr [3:5]) {} - #pragma omp target teams map(r.ArrS[3:5].A) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target teams map(r.ArrS [3:5].A) // expected-error {{OpenMP array section is not allowed here}} {} - #pragma omp target teams map(r.ArrS[3:5].Arr[6:7]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target teams map(r.ArrS [3:5].Arr [6:7]) // expected-error {{OpenMP array section is not allowed here}} {} - #pragma omp target teams map(r.ArrS[3].Arr[6:7]) +#pragma omp target teams map(r.ArrS[3].Arr [6:7]) {} - #pragma omp target teams map(r.S.Arr[4:5]) +#pragma omp target teams map(r.S.Arr [4:5]) {} - #pragma omp target teams map(r.S.Ptr[4:5]) +#pragma omp target teams map(r.S.Ptr [4:5]) {} - #pragma omp target teams map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target teams map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}} - #pragma omp target teams map((p+1)->A) +#pragma omp target teams map((p + 1)->A) {} - #pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}} +#pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}} {} - #pragma omp target data map(to: r.C) //expected-note {{used here}} +#pragma omp target data map(to \ + : r.C) //expected-note {{used here}} { - #pragma omp target teams map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} +#pragma omp target teams map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} {} } - #pragma omp target data map(to: t.Ptr) //expected-note {{used here}} +#pragma omp target data map(to \ + : t.Ptr) //expected-note {{used here}} { - #pragma omp target teams map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} {} } - #pragma omp target data map(to: t.C, t.D) +#pragma omp target data map(to \ + : t.C, t.D) { - #pragma omp target data map(to: t.C) +#pragma omp target data map(to \ + : t.C) { - #pragma omp target teams map(t.D) +#pragma omp target teams map(t.D) {} } } - #pragma omp target data map(to: t) +#pragma omp target data map(to \ + : t) { - #pragma omp target data map(to: t.C) +#pragma omp target data map(to \ + : t.C) { - #pragma omp target teams map(t.D) +#pragma omp target teams map(t.D) {} } } @@ -381,6 +426,16 @@ virtual void foo(); }; +class S9 { +public: + virtual void foo() {} +} s9; + +class S10 : public S9 { + int a; + void foo(){}; +} s10; + S3 h; #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}} @@ -443,9 +498,11 @@ foo(); #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target data map(tofrom: argc > 0 ? x : y) +// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target data map(tofrom \ + : x ?: y) #pragma omp target data map(argc) #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}} @@ -526,9 +583,11 @@ foo(); #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -// ge50-error@+2 {{expected addressable lvalue in 'map' clause}} // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) +// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target data map(tofrom \ + : argv[1] ?: argv[2]) #pragma omp target data map(argc) #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}} #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} @@ -613,6 +672,24 @@ {} #pragma omp target teams map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}} {} +#pragma omp target teams map(reinterpret_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*reinterpret_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map(static_cast (s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*static_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map(*dynamic_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map(const_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*const_cast (&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} +#pragma omp target teams map((S9)s10) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}} + {} +#pragma omp target teams map(*((S9 *)&s10)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} #pragma omp target teams map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams'}} {} Index: clang/test/OpenMP/target_update_codegen.cpp =================================================================== --- clang/test/OpenMP/target_update_codegen.cpp +++ clang/test/OpenMP/target_update_codegen.cpp @@ -1645,6 +1645,170 @@ #pragma omp target update to(arr[0:2][arg:][1:4], x, farr[0:2][1:2]) { ++arg; } } +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK26 -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 CK26 --check-prefix CK26-64 +// RUN: %clang_cc1 -DCK26 -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 CK26 --check-prefix CK26-64 +// RUN: %clang_cc1 -DCK26 -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 CK26 --check-prefix CK26-32 +// RUN: %clang_cc1 -DCK26 -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 CK26 --check-prefix CK26-32 + +// RUN: %clang_cc1 -DCK26 -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 -DCK26 -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 -DCK26 -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 -DCK26 -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 CK26 + +// CK26: id +int &id(int &x) { + return x; +} + +// CK26: foo +void foo(int arg) { + int x; + + // CK26: %x = alloca i32, align {{4|8}} + // CK26: [[OFFLOAD_BP:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK26: [[OFFLOAD_PTRS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK26: [[OFFLOAD_MAPPERS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK26: [[RESULT:%.+]] = call {{.+}} i32* @_Z2idRi(i32* {{.+}} %x) + // CK26: [[R0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK26: [[R1:%.+]] = bitcast i8** [[R0]] to i32* (i32*)** + // CK26: store i32* (i32*)* @_Z2idRi, i32* (i32*)** [[R1]], align {{4|8}} + // CK26: [[R2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK26: [[R3:%.+]] = bitcast i8** [[R2]] to i32** + // CK26: store i32* [[RESULT]], i32** [[R3]], align {{4|8}} + // CK26: [[R4:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_MAPPERS]], {{i32|i64}} 0, {{i32|i64}} 0 + // CK26: store i8* null, i8** [[R4]], align {{4|8}} + // CK26: [[GEPBP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK26: [[GEPP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK26-DAG: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** [[GEPBP]], i8** [[GEPP]], {{.+}}getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) + +#pragma omp target update to(id(x)) +} +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK27 -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 CK27 --check-prefix CK27-64 +// RUN: %clang_cc1 -DCK27 -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 CK27 --check-prefix CK27-64 +// RUN: %clang_cc1 -DCK27 -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 CK27 --check-prefix CK27-32 +// RUN: %clang_cc1 -DCK27 -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 CK27 --check-prefix CK27-32 + +// RUN: %clang_cc1 -DCK27 -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 -DCK27 -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 -DCK27 -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 -DCK27 -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 CK27 + +// CK27: wrapper +struct wrapper { + int wrapped; +}; + +// CK27: foo +void foo(int arg) { + wrapper x; + x.wrapped = 1; + wrapper *xp = &x; + + // CK27: %x = alloca %struct.wrapper, align 4 + // CK27: %xp = alloca %struct.wrapper*, align {{4|8}} + // CK27: [[OFFLOAD_BP:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK27: [[OFFLOAD_PTRS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK27: [[OFFLOAD_MAPPERS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK27: store i32 %arg, i32* %arg.addr, align 4 + // CK27: %wrapped = getelementptr inbounds %struct.wrapper, %struct.wrapper* %x, i32 0, i32 0 + // CK27: store i32 1, i32* %wrapped, align 4 + // CK27: store %struct.wrapper* %x, %struct.wrapper** %xp, align {{4|8}} + // CK27: [[R0:%.+]] = load %struct.wrapper*, %struct.wrapper** %xp, align {{4|8}} + // CK27: [[R1:%.+]] = load %struct.wrapper*, %struct.wrapper** %xp, align {{4|8}} + // CK27: [[R2:%.+]] = bitcast %struct.wrapper* [[R1]] to i32* + // CK27: [[R3:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK27: [[R4:%.+]] = bitcast i8** [[R3]] to %struct.wrapper** + // CK27: store %struct.wrapper* [[R0]], %struct.wrapper** [[R4]], align {{4|8}} + // CK27: [[R5:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK27: [[R6:%.+]] = bitcast i8** [[R5]] to i32** + // CK27: store i32* [[R2]], i32** [[R6]], align {{4|8}} + // CK27: [[R7:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_MAPPERS]], {{i32|i64}} 0, {{i32|i64}} 0 + // CK27: store i8* null, i8** [[R7]], align {{4|8}} + // CK27: [[GEPBP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK27: [[GEPP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK27: call void @__tgt_target_data_update_mapper(%struct.ident_t* @1, i64 -1, i32 1, i8** [[GEPBP]], i8** [[GEPP]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +#pragma omp target update to(*reinterpret_cast (xp)) +} +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK28 -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 CK28 --check-prefix CK28-64 +// RUN: %clang_cc1 -DCK28 -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 CK28 --check-prefix CK28-64 +// RUN: %clang_cc1 -DCK28 -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 CK28 --check-prefix CK28-32 +// RUN: %clang_cc1 -DCK28 -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 CK28 --check-prefix CK28-32 + +// RUN: %clang_cc1 -DCK28 -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 -DCK28 -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 -DCK28 -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 -DCK28 -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 CK28 + +// CK28: foo +void foo(int arg) { + bool which = arg > 0; + int a = 1; + int b = 2; + + // CK28: %arg.addr = alloca i32, align 4 + // CK28: %which = alloca i8, align 1 + // CK28: %a = alloca i32, align 4 + // CK28: %b = alloca i32, align 4 + // CK28: [[OFFLOAD_BP:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK28: [[OFFLOAD_PTRS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK28: [[OFFLOAD_MAPPERS:%.+]] = alloca [1 x i8*], align {{4|8}} + // CK28: store i32 %arg, i32* %arg.addr, align 4 + // CK28: [[R0:%.+]] = load i32, i32* %arg.addr, align 4 + // CK28: %cmp = icmp sgt i32 [[R0]], 0 + // CK28: %frombool = zext i1 %cmp to i8 + // CK28: store i8 %frombool, i8* %which, align 1 + // CK28: store i32 1, i32* %a, align 4 + // CK28: store i32 2, i32* %b, align 4 + // CK28: [[R1:%.+]] = load i8, i8* %which, align 1 + // CK28: %tobool = trunc i8 [[R1]] to i1 + // CK28: br i1 %tobool, label %[[TRUE_BRANCH:.+]], label %[[FALSE_BRANCH:.+]] + + // CK28: [[TRUE_BRANCH]]: ; preds = %entry + // CK28: br label %[[COND_END:.+]] + + // CK28: [[FALSE_BRANCH]]: ; preds = %entry + // CK28: br label %[[COND_END]] + + // CK28: [[COND_END]]: ; preds = %[[FALSE_BRANCH]], %[[TRUE_BRANCH]] + // CK28: [[LVALUE:%.+]] = phi i32* [ %a, %[[TRUE_BRANCH]] ], [ %b, %[[FALSE_BRANCH]] ] + // CK28: [[R2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK28: [[R3:%.+]] = bitcast i8** [[R2]] to i32** + // CK28: store i32* %a, i32** [[R3]], align {{4|8}} + // CK28: [[R4:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK28: [[R5:%.+]] = bitcast i8** [[R4]] to i32** + // CK28: store i32* [[LVALUE]], i32** [[R5]], align {{4|8}} + // CK28: [[R6:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_MAPPERS]], {{i32|i64}} 0, {{i32|i64}} 0 + // CK28: store i8* null, i8** [[R6]], align {{4|8}} + // CK28: [[GEPBP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP]], i32 0, i32 0 + // CK28: [[GEPP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 + // CK28: call void @__tgt_target_data_update_mapper(%struct.ident_t* @1, i64 -1, i32 1, i8** [[GEPBP]], i8** [[GEPP]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) + +#pragma omp target update to(which ? a : b) +} #endif #endif Index: clang/test/OpenMP/target_update_from_messages.cpp =================================================================== --- clang/test/OpenMP/target_update_from_messages.cpp +++ clang/test/OpenMP/target_update_from_messages.cpp @@ -1,9 +1,9 @@ -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized void foo() { } @@ -12,6 +12,14 @@ return argc; } +int &id(int &x) { + return x; +} + +int *id2(int *x) { + return x; +} + struct S1; // expected-note 2 {{declared here}} extern S1 a; class S2 { @@ -75,6 +83,12 @@ #pragma omp target update from(*(a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(*(*(this->ptr)+a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(*(this+this)) // expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}} + +#pragma omp target update from(id) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id(this->a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id(id(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*id2(&(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id2(&(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} } }; @@ -112,7 +126,8 @@ #pragma omp target update from(S2::S2sc) #pragma omp target update from(from) #pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}} -#pragma omp target update from(argc > 0 ? x : y) // le50-error 2 {{expected addressable lvalue in 'from' clause}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(argc > 0 ? x : y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(x ?: y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} @@ -141,6 +156,16 @@ return 0; } +class S9 { +public: + virtual void foo() {} +} s9; + +class S10 : public S9 { + int a; + void foo() {} +} s10; + int main(int argc, char **argv) { const int d = 5; const int da[5] = { 0 }; @@ -169,7 +194,8 @@ #pragma omp target update from(S2::S2sc) #pragma omp target update from(from) #pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}} -#pragma omp target update from(argc > 0 ? x : y) // le50-error {{expected addressable lvalue in 'from' clause}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(argc > 0 ? x : y) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(x ?: y) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} #pragma omp target update from(ba) @@ -204,6 +230,21 @@ { #pragma omp target update from(s7.x) } +#pragma omp target update from(id) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id(x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id(id(x))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*id2(&x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(id2(&x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + +#pragma omp target update from(reinterpret_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*reinterpret_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(static_cast (s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*static_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*dynamic_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(const_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*const_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from((S9)s10) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'from' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(*((S9 *)&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} return tmain(argc)+tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} expected-note {{in instantiation of function template specialization 'tmain' requested here}} } Index: clang/test/OpenMP/target_update_to_messages.cpp =================================================================== --- clang/test/OpenMP/target_update_to_messages.cpp +++ clang/test/OpenMP/target_update_to_messages.cpp @@ -1,12 +1,12 @@ -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized void foo() { } @@ -15,6 +15,14 @@ return argc; } +int &id(int &x) { + return x; +} + +int *id2(int *x) { + return x; +} + struct S1; // expected-note 2 {{declared here}} extern S1 a; class S2 { @@ -83,6 +91,12 @@ double marr[10][5][10]; #pragma omp target update to(marr [0:1][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} {} + +#pragma omp target update to(id) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id(this->a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id(id(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*id2(&(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id2(&(this->a))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} } }; @@ -120,7 +134,8 @@ #pragma omp target update to(S2::S2sc) #pragma omp target update to(to) #pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}} -#pragma omp target update to(argc > 0 ? x : y) // le50-error 2 {{expected addressable lvalue in 'to' clause}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(argc > 0 ? x : y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(x ?: y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} #pragma omp target update to(ba) @@ -147,6 +162,16 @@ return 0; } +class S9 { +public: + virtual void foo() {} +} s9; + +class S10 : public S9 { + int a; + void foo() {} +} s10; + int main(int argc, char **argv) { const int d = 5; const int da[5] = { 0 }; @@ -175,7 +200,8 @@ #pragma omp target update to(S2::S2sc) #pragma omp target update to(to) #pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}} -#pragma omp target update to(argc > 0 ? x : y) // le50-error {{expected addressable lvalue in 'to' clause}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(argc > 0 ? x : y) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(x ?: y) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} #pragma omp target update to(ba) @@ -212,6 +238,23 @@ #pragma omp target update to(*(m+(m))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}} #pragma omp target update to(*(1+y+y)) // expected-error {{indirection requires pointer operand ('int' invalid)}} {} + +#pragma omp target update to(id) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id(x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id(id(x))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*id2(&x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(id2(&x)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + +#pragma omp target update to(reinterpret_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*reinterpret_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(static_cast (s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*static_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*dynamic_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(const_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*const_cast (&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to((S9)s10) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected addressable lvalue in 'to' clause}} le50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(*((S9 *)&s10)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + return tmain(argc)+tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} expected-note {{in instantiation of function template specialization 'tmain' requested here}} }