Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7714,6 +7714,24 @@ "expected variable name%select{| or data member of current class}0">; def err_omp_expected_var_name_member_expr_or_array_item : Error< "expected variable name%select{|, data member of current class}0, array element or array section">; +def err_omp_expected_named_var_member_or_array_expression: Error< + "expected expression containing only member accesses and/or array sections based on named variables">; +def err_omp_bit_fields_forbidden_in_map_clause : Error< + "bit fields cannot be used to specify storage in a map clause">; +def err_omp_array_section_in_rightmost_expression : Error< + "array section can only be associated with the rightmost variable in a map clause expression">; +def err_omp_union_type_not_allowed : Error< + "mapped storage cannot be derived from a union">; +def err_omp_expected_access_to_data_field : Error< + "expected access to data field">; +def err_omp_multiple_array_items_in_map_clause : Error< + "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">; +def err_omp_pointer_mapped_along_with_derived_section : Error< + "pointer cannot be mapped along with a section derived from itself">; +def err_omp_original_storage_is_shared_and_does_not_contain : Error< + "original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage">; +def err_omp_same_pointer_derreferenced : Error< + "same pointer derreferenced in multiple different ways in map clause expressions">; def note_omp_task_predetermined_firstprivate_here : Note< "predetermined as a firstprivate in a task construct here">; def err_omp_threadprivate_incomplete_type : Error< Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -900,7 +900,6 @@ OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; bool MapTypeIsImplicit = false; bool MapTypeModifierSpecified = false; - bool UnexpectedId = false; SourceLocation DepLinMapLoc; // Parse '('. @@ -1021,7 +1020,8 @@ MapTypeIsImplicit = true; } } else { - UnexpectedId = true; + MapType = OMPC_MAP_tofrom; + MapTypeIsImplicit = true; } if (Tok.is(tok::colon)) { @@ -1036,7 +1036,7 @@ ((Kind != OMPC_reduction) && (Kind != OMPC_depend) && (Kind != OMPC_map)) || ((Kind == OMPC_reduction) && !InvalidReductionId) || - ((Kind == OMPC_map) && (UnexpectedId || MapType != OMPC_MAP_unknown) && + ((Kind == OMPC_map) && (MapType != OMPC_MAP_unknown) && (!MapTypeModifierSpecified || (MapTypeModifierSpecified && MapTypeModifier == OMPC_MAP_always))) || ((Kind == OMPC_depend) && DepKind != OMPC_DEPEND_unknown); Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -77,12 +77,9 @@ ImplicitDSALoc() {} }; -public: - struct MapInfo { - Expr *RefExpr; - }; - private: + typedef SmallVector MapInfo; + struct DSAInfo { OpenMPClauseKind Attributes; Expr *RefExpr; @@ -337,30 +334,39 @@ Scope *getCurScope() { return Stack.back().CurScope; } SourceLocation getConstructLoc() { return Stack.back().ConstructLoc; } - MapInfo getMapInfoForVar(ValueDecl *VD) { - MapInfo VarMI = {0}; - for (auto Cnt = Stack.size() - 1; Cnt > 0; --Cnt) { - if (Stack[Cnt].MappedDecls.count(VD)) { - VarMI = Stack[Cnt].MappedDecls[VD]; - break; - } + // Do the check specified in MapInfoCheck and return true if any issue is + // found. + template + bool checkMapInfoForVar(ValueDecl *VD, bool CurrentRegionOnly, + MapInfoCheck Check) { + auto SI = Stack.rbegin(); + auto SE = Stack.rend(); + + if (SI == SE) + return false; + + if (CurrentRegionOnly) { + SE = std::next(SI); + } else { + ++SI; } - return VarMI; - } - void addMapInfoForVar(ValueDecl *VD, MapInfo MI) { - if (Stack.size() > 1) { - Stack.back().MappedDecls[VD] = MI; + for (; SI != SE; ++SI) { + auto MI = SI->MappedDecls.find(VD); + if (MI != SI->MappedDecls.end()) { + for (Expr *E : MI->second) { + if (Check(E)) + return true; + } + } } + return false; } - MapInfo IsMappedInCurrentRegion(ValueDecl *VD) { - assert(Stack.size() > 1 && "Target level is 0"); - MapInfo VarMI = {0}; - if (Stack.size() > 1 && Stack.back().MappedDecls.count(VD)) { - VarMI = Stack.back().MappedDecls[VD]; + void addExprToVarMapInfo(ValueDecl *VD, Expr *E) { + if (Stack.size() > 1) { + Stack.back().MappedDecls[VD].push_back(E); } - return VarMI; } }; bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) { @@ -8555,6 +8561,387 @@ return true; } +// Return the expression of the base of the map clause or null if it cannot +// be determined and do all the necessary checks to see if the expression is +// valid as a standalone map clause expression. +static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) { + SourceLocation ELoc = E->getExprLoc(); + SourceRange ERange = E->getSourceRange(); + + // The base of elements of list in a map clause have to be either: + // - a reference to variable or field. + // - a member expression. + // - an array expression. + // + // E.g. if we have the expression 'r.S.Arr[:12]', we want to retrieve the + // reference to 'r'. + // + // If we have: + // + // struct SS { + // Bla S; + // foo() { + // #pragma omp target map (S.Arr[:12]); + // } + // } + // + // We want to retrieve the member expression 'this->S'; + + Expr *RelevantExpr = nullptr; + + // Flags to help capture some memory + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2] + // If a list item is an array section, it must specify contiguous storage. + // + // For this restriction it is sufficient that we make sure only references + // to variables or fields and array expressions, and that no array sections + // exist except in the rightmost expression. E.g. these would be invalid: + // + // r.ArrS[3:5].Arr[6:7] + // + // r.ArrS[3:5].x + // + // but these would be valid: + // r.ArrS[3].Arr[6:7] + // + // r.ArrS[3].x + + bool IsRightMostExpression = true; + + while (!RelevantExpr) { + auto AllowArraySection = IsRightMostExpression; + IsRightMostExpression = false; + + E = E->IgnoreParenImpCasts(); + + if (auto *CurE = dyn_cast(E)) { + if (!isa(CurE->getDecl())) + break; + + RelevantExpr = CurE; + continue; + } + + if (auto *CurE = dyn_cast(E)) { + auto *BaseE = CurE->getBase()->IgnoreParenImpCasts(); + + if (isa(BaseE)) + // We found a base expression: this->Val. + RelevantExpr = CurE; + else + E = BaseE; + + if (!isa(CurE->getMemberDecl())) { + SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field) + << CurE->getSourceRange(); + break; + } + + auto *FD = cast(CurE->getMemberDecl()); + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3] + // A bit-field cannot appear in a map clause. + // + if (FD->isBitField()) { + SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_map_clause) + << CurE->getSourceRange(); + break; + } + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1] + // If the type of a list item is a reference to a type T then the type + // will be considered to be T for all purposes of this clause. + QualType CurType = BaseE->getType(); + if (CurType->isReferenceType()) + CurType = CurType->getPointeeType(); + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.2] + // A list item cannot be a variable that is a member of a structure with + // a union type. + // + if (auto *RT = CurType->getAs()) + if (RT->isUnionType()) { + SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed) + << CurE->getSourceRange(); + break; + } + + continue; + } + + if (auto *CurE = dyn_cast(E)) { + E = CurE->getBase()->IgnoreParenImpCasts(); + + if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) { + SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) + << 0 << CurE->getSourceRange(); + break; + } + continue; + } + + if (auto *CurE = dyn_cast(E)) { + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7] + // If a list item is an element of a structure, only the rightmost symbol + // of the variable reference can be an array section. + // + if (!AllowArraySection) { + SemaRef.Diag(ELoc, diag::err_omp_array_section_in_rightmost_expression) + << CurE->getSourceRange(); + break; + } + + E = CurE->getBase()->IgnoreParenImpCasts(); + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1] + // If the type of a list item is a reference to a type T then the type + // will be considered to be T for all purposes of this clause. + QualType CurType = E->getType(); + if (CurType->isReferenceType()) + CurType = CurType->getPointeeType(); + + if (!CurType->isAnyPointerType() && !CurType->isArrayType()) { + SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) + << 0 << CurE->getSourceRange(); + break; + } + + continue; + } + + // If nothing else worked, this is not a valid map clause expression. + SemaRef.Diag(ELoc, + diag::err_omp_expected_named_var_member_or_array_expression) + << ERange; + break; + } + + return RelevantExpr; +} + +// Return true if expression E associated with value VD has conflicts with other +// map information. +static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD, + Expr *E, bool CurrentRegionOnly) { + assert(VD && E); + + // Types used to organize the components of a valid map clause. + typedef std::pair MapExpressionComponent; + typedef SmallVector MapExpressionComponents; + + // Helper to extract the components in the map clause expression E and store + // them into MEC. This assumes that E is a valid map clause expression, i.e. + // it has already passed the single clause checks. + auto ExtractMapExpressionComponents = [](Expr *TE, + MapExpressionComponents &MEC) { + while (true) { + TE = TE->IgnoreParenImpCasts(); + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back( + MapExpressionComponent(CurE, cast(CurE->getDecl()))); + break; + } + + if (auto *CurE = dyn_cast(TE)) { + auto *BaseE = CurE->getBase()->IgnoreParenImpCasts(); + + MEC.push_back(MapExpressionComponent( + CurE, cast(CurE->getMemberDecl()))); + if (isa(BaseE)) + break; + + TE = BaseE; + continue; + } + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back(MapExpressionComponent(CurE, nullptr)); + TE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back(MapExpressionComponent(CurE, nullptr)); + TE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + llvm_unreachable( + "Expecting only valid map clause expressions at this point!"); + } + }; + + SourceLocation ELoc = E->getExprLoc(); + SourceRange ERange = E->getSourceRange(); + + // In order to easily check the conflicts we need to match each component of + // the expression under test with the components of the expressions that are + // already in the stack. + + MapExpressionComponents CurComponents; + ExtractMapExpressionComponents(E, CurComponents); + + assert(!CurComponents.empty() && "Map clause expression with no components!"); + assert(CurComponents.back().second == VD && + "Map clause expression with unexpected base!"); + + // Variables to help detecting enclosing problems in data environment nests. + bool IsEnclosedByDataEnvironmentExpr = false; + Expr *EnclosingExpr = nullptr; + + bool FoundError = + DSAS->checkMapInfoForVar(VD, CurrentRegionOnly, [&](Expr *RE) -> bool { + MapExpressionComponents StackComponents; + ExtractMapExpressionComponents(RE, StackComponents); + assert(!StackComponents.empty() && + "Map clause expression with no components!"); + assert(StackComponents.back().second == VD && + "Map clause expression with unexpected base!"); + + // Expressions must start from the same base. Here we detect at which + // point both expressions diverge from each other and see if we can + // detect if the memory referred to both expressions is contiguous and + // do not overlap. + auto CI = CurComponents.rbegin(); + auto CE = CurComponents.rend(); + auto SI = StackComponents.rbegin(); + auto SE = StackComponents.rend(); + for (; CI != CE && SI != SE; ++CI, ++SI) { + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.3] + // At most one list item can be an array item derived from a given + // variable in map clauses of the same construct. + if (CurrentRegionOnly && (isa(CI->first) || + isa(CI->first)) && + (isa(SI->first) || + isa(SI->first))) { + SemaRef.Diag(CI->first->getExprLoc(), + diag::err_omp_multiple_array_items_in_map_clause) + << CI->first->getSourceRange(); + ; + SemaRef.Diag(SI->first->getExprLoc(), diag::note_used_here) + << SI->first->getSourceRange(); + return true; + } + + // Do both expressions have the same kind? + if (CI->first->getStmtClass() != SI->first->getStmtClass()) + break; + + // Are we dealing with different variables/fields? + if (CI->second != SI->second) + break; + } + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4] + // List items of map clauses in the same construct must not share + // original storage. + // + // If the expressions are exactly the same or one is a subset of the + // other, it means they are sharing storage. + if (CI == CE && SI == SE) { + if (CurrentRegionOnly) { + SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + SemaRef.Diag(RE->getExprLoc(), diag::note_used_here) + << RE->getSourceRange(); + return true; + } else { + // If we find the same expression in the enclosing data environment, + // that is legal. + IsEnclosedByDataEnvironmentExpr = true; + return false; + } + } + + QualType DerivedType = std::prev(CI)->first->getType(); + SourceLocation DerivedLoc = std::prev(CI)->first->getExprLoc(); + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1] + // If the type of a list item is a reference to a type T then the type + // will be considered to be T for all purposes of this clause. + if (DerivedType->isReferenceType()) + DerivedType = DerivedType->getPointeeType(); + + // 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 + // clauses of the same construct. + // + // Also, cover one of the cases in: + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.5] + // If any part of the original storage of a list item has corresponding + // storage in the device data environment, all of the original storage + // must have corresponding storage in the device data environment. + // + if (DerivedType->isAnyPointerType()) { + if (CI == CE || SI == SE) { + SemaRef.Diag( + DerivedLoc, + diag::err_omp_pointer_mapped_along_with_derived_section) + << DerivedLoc; + } else { + assert(CI != CE && SI != SE); + SemaRef.Diag(DerivedLoc, diag::err_omp_same_pointer_derreferenced) + << DerivedLoc; + } + SemaRef.Diag(RE->getExprLoc(), diag::note_used_here) + << RE->getSourceRange(); + return true; + } + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4] + // List items of map clauses in the same construct must not share + // original storage. + // + // An expression is a subset of the other. + if (CurrentRegionOnly && (CI == CE || SI == SE)) { + SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + SemaRef.Diag(RE->getExprLoc(), diag::note_used_here) + << RE->getSourceRange(); + return true; + } + + // The current expression uses the same base as other expression in the + // data environment but does not contain it completelly. + if (!CurrentRegionOnly && SI != SE) + EnclosingExpr = RE; + + // The current expression is a subset of the expression in the data + // environment. + IsEnclosedByDataEnvironmentExpr |= + (!CurrentRegionOnly && CI != CE && SI == SE); + + return false; + }); + + if (CurrentRegionOnly) + return FoundError; + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.5] + // If any part of the original storage of a list item has corresponding + // storage in the device data environment, all of the original storage must + // have corresponding storage in the device data environment. + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.6] + // If a list item is an element of a structure, and a different element of + // the structure has a corresponding list item in the device data environment + // prior to a task encountering the construct associated with the map clause, + // then the list item must also have a correspnding list item in the device + // data environment prior to the task encountering the construct. + // + if (EnclosingExpr && !IsEnclosedByDataEnvironmentExpr) { + SemaRef.Diag(ELoc, + diag::err_omp_original_storage_is_shared_and_does_not_contain) + << ERange; + SemaRef.Diag(EnclosingExpr->getExprLoc(), diag::note_used_here) + << EnclosingExpr->getSourceRange(); + return true; + } + + return FoundError; +} + OMPClause * Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, @@ -8572,96 +8959,82 @@ } SourceLocation ELoc = RE->getExprLoc(); - // OpenMP [2.14.5, Restrictions] - // A variable that is part of another variable (such as field of a - // structure) but is not an array element or an array section cannot appear - // in a map clause. auto *VE = RE->IgnoreParenLValueCasts(); if (VE->isValueDependent() || VE->isTypeDependent() || VE->isInstantiationDependent() || VE->containsUnexpandedParameterPack()) { - // It will be analyzed later. + // We can only analyze this information once the missing information is + // resolved. Vars.push_back(RE); continue; } auto *SimpleExpr = RE->IgnoreParenCasts(); - auto *DE = dyn_cast(SimpleExpr); - auto *ASE = dyn_cast(SimpleExpr); - auto *OASE = dyn_cast(SimpleExpr); - - if (!RE->IgnoreParenImpCasts()->isLValue() || - (!OASE && !ASE && !DE) || - (DE && !isa(DE->getDecl())) || - (ASE && !ASE->getBase()->getType()->isAnyPointerType() && - !ASE->getBase()->getType()->isArrayType())) { - Diag(ELoc, diag::err_omp_expected_var_name_member_expr_or_array_item) - << 0 << RE->getSourceRange(); + + if (!RE->IgnoreParenImpCasts()->isLValue()) { + Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression) + << RE->getSourceRange(); continue; } - Decl *D = nullptr; - if (DE) { - D = DE->getDecl(); - } else if (ASE) { - auto *B = ASE->getBase()->IgnoreParenCasts(); - D = dyn_cast(B)->getDecl(); - } else if (OASE) { - auto *B = OASE->getBase(); - D = dyn_cast(B)->getDecl(); + // Obtain the array or member expression bases if required. + auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr); + if (!BE) + continue; + + // If the base is a reference to a variable, we rely on that variable for + // the following checks. If it is a 'this' expression we rely on the field. + ValueDecl *D = nullptr; + if (auto *DRE = dyn_cast(BE)) { + D = DRE->getDecl(); + } else { + auto *ME = cast(BE); + assert(isa(ME->getBase()) && "Unexpected expression!"); + D = ME->getMemberDecl(); } assert(D && "Null decl on map clause."); - auto *VD = cast(D); - // OpenMP [2.14.5, Restrictions, p.8] - // threadprivate variables cannot appear in a map clause. - if (DSAStack->isThreadPrivate(VD)) { + auto *VD = dyn_cast(D); + auto *FD = dyn_cast(D); + + assert((VD || FD) && "Only variables or fields are expected here!"); + + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10] + // threadprivate variables cannot appear in a map clause. + if (VD && DSAStack->isThreadPrivate(VD)) { auto DVar = DSAStack->getTopDSA(VD, false); Diag(ELoc, diag::err_omp_threadprivate_in_map); ReportOriginalDSA(*this, DSAStack, VD, DVar); continue; } - // OpenMP [2.14.5, Restrictions, p.2] - // At most one list item can be an array item derived from a given variable - // in map clauses of the same construct. - // OpenMP [2.14.5, Restrictions, p.3] - // List items of map clauses in the same construct must not share original - // storage. - // OpenMP [2.14.5, Restrictions, C/C++, p.2] - // A variable for which the type is pointer, reference to array, or - // reference to pointer and an array section derived from that variable - // must not appear as list items of map clauses of the same construct. - DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD); - if (MI.RefExpr) { - Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc; - Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) - << MI.RefExpr->getSourceRange(); - continue; - } + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9] + // A list item cannot appear in both a map clause and a data-sharing + // attribute clause on the same construct. + // + // TODO: Implement this check - it cannot currently be tested because of + // missing implementation of the other data sharing clauses in target + // directives. + + // Check conflicts with other map clause expressions. We check the conflicts + // with the current construct separately from the enclosing data + // environment, because the restrictions are different. + if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr, + /*CurrentRegionOnly=*/true)) + break; + if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr, + /*CurrentRegionOnly=*/false)) + break; - // OpenMP [2.14.5, Restrictions, C/C++, p.3,4] - // A variable for which the type is pointer, reference to array, or - // reference to pointer must not appear as a list item if the enclosing - // device data environment already contains an array section derived from - // that variable. - // An array section derived from a variable for which the type is pointer, - // reference to array, or reference to pointer must not appear as a list - // item if the enclosing device data environment already contains that - // variable. - QualType Type = VD->getType(); - MI = DSAStack->getMapInfoForVar(VD); - if (MI.RefExpr && (isa(MI.RefExpr->IgnoreParenLValueCasts()) != - isa(VE)) && - (Type->isPointerType() || Type->isReferenceType())) { - Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc; - Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) - << MI.RefExpr->getSourceRange(); - continue; - } + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1] + // If the type of a list item is a reference to a type T then the type will + // be considered to be T for all purposes of this clause. + QualType Type = D->getType(); + if (Type->isReferenceType()) + Type = Type->getPointeeType(); - // OpenMP [2.14.5, Restrictions, C/C++, p.7] + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9] // A list item must have a mappable type. if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this, DSAStack, Type)) @@ -8699,8 +9072,7 @@ } Vars.push_back(RE); - MI.RefExpr = RE; - DSAStack->addMapInfoForVar(VD, MI); + DSAStack->addExprToVarMapInfo(D, RE); } if (Vars.empty()) return nullptr; Index: test/OpenMP/target_map_messages.cpp =================================================================== --- test/OpenMP/target_map_messages.cpp +++ test/OpenMP/target_map_messages.cpp @@ -1,5 +1,181 @@ // RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s +template +struct SA { + static int ss; + #pragma omp threadprivate(ss) // expected-note {{defined as threadprivate or thread local}} + float a; + int b[12]; + float *c; + T d; + float e[I]; + T *f; + void func(int arg) { + #pragma omp target map(arg,a,d) + {} + #pragma omp target map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}} + {} + #pragma omp target map(arg,a*2) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} + #pragma omp target map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} + #pragma omp target map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}} + {} + #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}} + {} + + #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) 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,e[:]) + {} + + #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,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}} + {} + return; + } +}; + +struct SB { + unsigned A; + unsigned B; + float Arr[100]; + float *Ptr; + float *foo() { + return &Arr[0]; + } +}; + +struct SC { + unsigned A : 2; + unsigned B : 3; + unsigned C; + unsigned D; + float Arr[100]; + SB S; + SB ArrS[100]; + SB *PtrS; + SB *&RPtrS; + float *Ptr; + + SC(SB *&_RPtrS) : RPtrS(_RPtrS) {} +}; + +union SD { + unsigned A; + float B; +}; + +void SAclient(int arg) { + SA s; + s.func(arg); // expected-note {{in instantiation of member function}} + + SB *p; + + SD u; + SC r(p),t(p); + #pragma omp target map(r) + {} + #pragma omp target map(r.ArrS[0].B) + {} + #pragma omp target map(r.ArrS[0].Arr[1:23]) + {} + #pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}} + {} + #pragma omp target 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 map(r.ArrS[0].A, t.ArrS[1].A) + {} + #pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}} + {} + #pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}} + {} + #pragma omp target map(r.S.Arr[:12]) + {} + #pragma omp target map(r.S.foo()[:12]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} + #pragma omp target map(r.C, r.D) + {} + #pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + {} + #pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + {} + #pragma omp target 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 map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + {} + #pragma omp target map(r.C, t.C) + {} + #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a map clause}} + {} + #pragma omp target map(r.Arr) + {} + #pragma omp target map(r.Arr[3:5]) + {} + #pragma omp target map(r.Ptr[3:5]) + {} + #pragma omp target map(r.ArrS[3:5].A) // expected-error {{OpenMP array section is not allowed here}} + {} + #pragma omp target map(r.ArrS[3:5].Arr[6:7]) // expected-error {{OpenMP array section is not allowed here}} + {} + #pragma omp target map(r.ArrS[3].Arr[6:7]) + {} + #pragma omp target map(r.S.Arr[4:5]) + {} + #pragma omp target map(r.S.Ptr[4:5]) + {} + #pragma omp target map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + {} + #pragma omp target map((p+1)->A) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + {} + #pragma omp target map(u.B) // expected-error {{mapped storage cannot be derived from a union}} + {} + + #pragma omp target data map(to: r.C) //expected-note {{used here}} + { + #pragma omp target 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 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) + { + #pragma omp target map(t.D) + {} + } + } + + #pragma omp target data map(to: t) + { + #pragma omp target data map(to: t.C) + { + #pragma omp target map(t.D) + {} + } + } +} + void foo() { } @@ -63,15 +239,22 @@ T to, tofrom, always; const T (&l)[5] = da; - #pragma omp target map // expected-error {{expected '(' after 'map'}} + {} #pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} + {} #pragma omp target map() // expected-error {{expected expression}} + {} #pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}} + {} #pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}} + {} #pragma omp target map(to:) // expected-error {{expected expression}} + {} #pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} #pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} + {} #pragma omp target map(x) foo(); #pragma omp target map(tofrom: t[:I]) @@ -80,7 +263,7 @@ foo(); #pragma omp target map(T) // expected-error {{'T' does not refer to a value}} foo(); -#pragma omp target map(I) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target map(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} foo(); #pragma omp target map(S2::S2s) foo(); @@ -96,39 +279,39 @@ foo(); #pragma omp target map(to, x) foo(); -#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -#pragma omp target map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} -#pragma omp target map(argc) -#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}} -#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}} -#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} -#pragma omp target map(ca) -#pragma omp target map(da) -#pragma omp target map(S2::S2s) -#pragma omp target map(S2::S2sc) -#pragma omp target map(e, g) -#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}} -#pragma omp target map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} -#pragma omp target map(k), map(k[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target data map(tofrom: argc > 0 ? x : y) // expected-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}} expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target data map(ca) +#pragma omp target data map(da) +#pragma omp target data map(S2::S2s) +#pragma omp target data map(S2::S2sc) +#pragma omp target data map(e, g) +#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target data map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} foo(); -#pragma omp target map(da) +#pragma omp target data map(da) #pragma omp target map(da[:4]) foo(); -#pragma omp target map(k, j, l) // expected-note 4 {{used here}} -#pragma omp target map(k[:4]) // expected-error 2 {{variable already marked as mapped in current construct}} -#pragma omp target map(j) -#pragma omp target map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} +#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} +#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(j) +#pragma omp target map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} foo(); -#pragma omp target map(k[:4], j, l[:5]) // expected-note 4 {{used here}} -#pragma omp target map(k) // expected-error 2 {{variable already marked as mapped in current construct}} -#pragma omp target map(j) -#pragma omp target map(l) // expected-error 2 {{variable already marked as mapped in current construct}} +#pragma omp target data map(k[:4], j, l[:5]) // expected-note 4 {{used here}} +#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(j) +#pragma omp target map(l) // expected-error 2 {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} foo(); -#pragma omp target map(always, tofrom: x) -#pragma omp target map(always: x) // expected-error {{missing map type}} -#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} -#pragma omp target map(always, tofrom: always, tofrom, x) +#pragma omp target data map(always, tofrom: x) +#pragma omp target data map(always: x) // expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target data map(always, tofrom: always, tofrom, x) #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}} foo(); @@ -147,14 +330,14 @@ int y; int to, tofrom, always; const int (&l)[5] = da; -#pragma omp target map // expected-error {{expected '(' after 'map'}} -#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} -#pragma omp target map() // expected-error {{expected expression}} -#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}} -#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}} -#pragma omp target map(to:) // expected-error {{expected expression}} -#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} -#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target data map // expected-error {{expected '(' after 'map'}} +#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} +#pragma omp target data map() // expected-error {{expected expression}} +#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}} +#pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target data map(to:) // expected-error {{expected expression}} +#pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} +#pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target map(x) foo(); #pragma omp target map(to: x) @@ -165,40 +348,40 @@ foo(); #pragma omp target map(to, x) foo(); -#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} -#pragma omp target map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}} -#pragma omp target map(argc) -#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}} -#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}} -#pragma omp target map(argv[1]) -#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} -#pragma omp target map(ca) -#pragma omp target map(da) -#pragma omp target map(S2::S2s) -#pragma omp target map(S2::S2sc) -#pragma omp target map(e, g) -#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}} -#pragma omp target map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} -#pragma omp target map(k), map(k[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#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]) // expected-error {{xpected 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}} expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target data map(argv[1]) +#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target data map(ca) +#pragma omp target data map(da) +#pragma omp target data map(S2::S2s) +#pragma omp target data map(S2::S2sc) +#pragma omp target data map(e, g) +#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} foo(); -#pragma omp target map(da) +#pragma omp target data map(da) #pragma omp target map(da[:4]) foo(); -#pragma omp target map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target map(k[:4]) // expected-error {{variable already marked as mapped in current construct}} -#pragma omp target map(j) -#pragma omp target map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} +#pragma omp target data map(k, j, l) // expected-note {{used here}} +#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(j) +#pragma omp target map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} foo(); -#pragma omp target map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -#pragma omp target map(k) // expected-error {{variable already marked as mapped in current construct}} -#pragma omp target map(j) -#pragma omp target map(l) // expected-error {{variable already marked as mapped in current construct}} +#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} +#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(j) +#pragma omp target map(l) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} foo(); -#pragma omp target map(always, tofrom: x) -#pragma omp target map(always: x) // expected-error {{missing map type}} -#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} -#pragma omp target map(always, tofrom: always, tofrom, x) +#pragma omp target data map(always, tofrom: x) +#pragma omp target data map(always: x) // expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target data map(always, tofrom: always, tofrom, x) #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}} foo();