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) @@ -17050,8 +17062,11 @@ } bool VisitUnaryOperator(UnaryOperator *UO) { - if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() || - UO->getOpcode() != UO_Deref) { + bool isFullExpression = Components.empty(); + + if (SemaRef.getLangOpts().OpenMP < 50 || + (isFullExpression && + (!UO->isLValue() || UO->getOpcode() != UO_Deref))) { emitErrorMsg(); return false; } @@ -17092,6 +17107,49 @@ 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 VisitConditionalOperator(ConditionalOperator *CO) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); + Components.emplace_back(CO, nullptr, false); + return RelevantExpr || Visit(CO->getTrueExpr()->IgnoreParenImpCasts()) || + Visit(CO->getFalseExpr()->IgnoreParenImpCasts()); + } + bool VisitBinaryConditionalOperator(BinaryConditionalOperator *BCO) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); + Components.emplace_back(BCO, nullptr, false); + Expr *TrueExpr = + dyn_cast(BCO->getTrueExpr())->getSourceExpr(); + assert(TrueExpr && "Failed to resolve true expression of Elvis operator."); + return RelevantExpr || Visit(TrueExpr->IgnoreParenImpCasts()) || + Visit(BCO->getFalseExpr()->IgnoreParenImpCasts()); + } bool VisitStmt(Stmt *) { emitErrorMsg(); return false; @@ -17268,7 +17326,7 @@ } QualType DerivedType = - std::prev(CI)->getAssociatedDeclaration()->getType(); + std::prev(CI)->getAssociatedExpression()->getType(); SourceLocation DerivedLoc = std::prev(CI)->getAssociatedExpression()->getExprLoc(); @@ -17277,6 +17335,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 +17360,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 +17388,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 +17424,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 +17726,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,23 +81,37 @@ {} #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}} @@ -166,11 +188,11 @@ // 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'}} @@ -181,9 +203,9 @@ // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} #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 +219,9 @@ float *foo() { return &Arr[0]; } + float &bar() { + return Arr[0]; + } }; struct SC { @@ -356,7 +381,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 +601,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 +702,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 +743,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 +877,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,6 +64,16 @@ {} #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}} {} @@ -97,6 +115,9 @@ float *foo() { return &Arr[0]; } + float &bar() { + return Arr[0]; + } }; struct SC { @@ -256,10 +277,17 @@ {} #pragma omp target teams map(r.S.Arr[:12]) {} - // ge50-error@+2 {{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}} 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.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} @@ -381,6 +409,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 +481,10 @@ 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 +565,10 @@ 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 +653,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}} }