diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9770,6 +9770,9 @@ def err_omp_lastprivate_conditional_non_scalar : Error< "expected list item of scalar type in 'lastprivate' clause with 'conditional' modifier" >; +def err_omp_non_lvalue_in_map_or_motion_clauses: Error< + "expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'" + >; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7406,6 +7406,7 @@ // S1 s; // double *p; // struct S2 *ps; + // struct S2 **ps_ptr; // } // S2 s; // S2 *ps; @@ -7542,6 +7543,48 @@ // &(ps->p), &(ps->p[0]), 33*sizeof(double), MEMBER_OF(4) | PTR_AND_OBJ | TO // (*) the struct this entry pertains to is the 4th element in the list // of arguments, hence MEMBER_OF(4) + // + // map(*p) + // &p, &(p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map(*(p+3)) + // &p, &(p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map((p+3)+7) + // &p, &(p[10]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map(*(ps->p)) + // ps, &(ps->p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map(*(ps->p+3)) + // ps, &(ps->p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map((ps->p+3)[1]) + // ps, &(ps->p[4]), 1*sizeof(float), TARGET_PARAM | TO | FROM + // + // map(*(ps->ps->p)) + // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM + // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1) + // &(ps->ps), &(ps->ps->p[0]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ | + // FROM + // + // map((ps->ps->p)[3]) + // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM + // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1) + // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ | + // FROM + // + // map(*(ps->ps->p+3)) + // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM + // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1) + // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ | + // FROM + // + // map(*(ps->ps->ps_ptr)) + // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM + // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1) + // &(ps->ps_ptr), &(ps->ps->ps_ptr[0]), sizeof(S2*), MEMBER_OF(1) | + // PTR_AND_OBJ | TO | FROM // Track if the map information being generated is the first for a capture. bool IsCaptureFirstInfo = IsFirstComponentList; @@ -7662,10 +7705,7 @@ if (Next == CE || IsPointer || IsFinalArraySection) { // If this is not the last component, we expect the pointer to be // associated with an array expression or member expression. - assert((Next == CE || - isa(Next->getAssociatedExpression()) || - isa(Next->getAssociatedExpression()) || - isa(Next->getAssociatedExpression())) && + assert((Next == CE || Next->getAssociatedExpression()->isLValue()) && "Unexpected expression"); Address LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()) @@ -7781,6 +7821,8 @@ // mapped member. If the parent is "*this", then the value declaration // is nullptr. if (EncounteredME) { + assert(!isa(EncounteredME->getMemberDecl()) && + "Unexpected FunctionDecl"); const auto *FD = dyn_cast(EncounteredME->getMemberDecl()); unsigned FieldIndex = FD->getFieldIndex(); diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -15197,256 +15197,411 @@ return ConstLength.getSExtValue() != 1; } -// Return the expression of the base of the mappable expression or null if it -// cannot be determined and do all the necessary checks to see if the expression -// is valid as a standalone mappable expression. In the process, record all the -// components of the expression. -static const Expr *checkMapClauseExpressionBase( - Sema &SemaRef, Expr *E, - OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose) { - 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'; - - const Expr *RelevantExpr = nullptr; - - // 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 (unless they cover the whole - // dimension of the array). 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 +namespace { +// 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'; - bool AllowUnitySizeArraySection = true; - bool AllowWholeSizeArraySection = true; +// 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 (unless they cover the whole +// dimension of the array). 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 +class MapBaseChecker final : public StmtVisitor { + Sema &SemaRef; + OpenMPClauseKind CKind; + OMPClauseMappableExprCommon::MappableExprComponentList &Components; + bool NoDiagnose; + const Expr *RelevantExpr{nullptr}; + size_t Depth{0}; + SmallVector CntStack; + bool AllowUnitySizeArraySection{true}; + bool AllowWholeSizeArraySection{true}; + SourceLocation ELoc; + SourceRange ERange; + + void emitErrorMsg() { + if (!NoDiagnose) { + // If nothing else worked, this is not a valid map clause expression. + if (SemaRef.getLangOpts().OpenMP < 50) { + SemaRef.Diag( + ELoc, diag::err_omp_expected_named_var_member_or_array_expression) + << ERange; + } else { + SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses) + << ERange; + } + } + } - while (!RelevantExpr) { - E = E->IgnoreParenImpCasts(); +public: + bool VisitDeclRefExpr(DeclRefExpr *DRE) { + if (!isa(DRE->getDecl())) + return false; - if (auto *CurE = dyn_cast(E)) { - if (!isa(CurE->getDecl())) - return nullptr; + if (SemaRef.getLangOpts().OpenMP < 50) { + RelevantExpr = DRE; + // Record the component. + Components.emplace_back(DRE, DRE->getDecl()); + } else { + const clang::Type *PT = DRE->getType().getTypePtr(); + const std::string TypeStr = PT->getCanonicalTypeInternal().getAsString(); + size_t TypeDecorCnt = std::count_if(TypeStr.begin(), TypeStr.end(), + [](char c) { return c == '*' || c == '['; }); + // Normally we want to add Components if Depth == TypeDecorCnt, for example: + // + // int **B, *l; + // map(**(B + *l)) + // Depth for B: 2 + // TypeDecorCnt for B: 2 + // Depth for l: 3 + // TypeDecorCnt for l: 1 + // + // However, combining with address_of operator(&), we might have Depth + // less than TypeDecorCnt as LValue + // + // int B; + // map(*(&B)) + // Depth for B: 1 + // TypeDecorCnt for B: 0 + if (Depth <= TypeDecorCnt) { + RelevantExpr = cast(DRE); + // record the component. + Components.emplace_back(DRE, DRE->getDecl()); + } + } - RelevantExpr = CurE; + return true; + } + bool VisitMemberExpr(MemberExpr *ME) { + Expr *E = cast(ME); + Expr *BaseE = ME->getBase()->IgnoreParenImpCasts(); + CntStack.emplace_back(Depth); + Depth = 0; - // If we got a reference to a declaration, we should not expect any array - // section before that. - AllowUnitySizeArraySection = false; - AllowWholeSizeArraySection = false; + if (isa(BaseE)) + // We found a base expression: this->Val. + RelevantExpr = ME; + else + E = BaseE; - // Record the component. - CurComponents.emplace_back(CurE, CurE->getDecl()); - } else if (auto *CurE = dyn_cast(E)) { - Expr *BaseE = CurE->getBase()->IgnoreParenImpCasts(); + if (!isa(ME->getMemberDecl())) { + if (!NoDiagnose) { + SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field) + << ME->getSourceRange(); + return false; + } + if (RelevantExpr) + return false; + return Visit(E->IgnoreParenImpCasts());; + } - if (isa(BaseE)) - // We found a base expression: this->Val. - RelevantExpr = CurE; - else - E = BaseE; + auto *FD = cast(ME->getMemberDecl()); - if (!isa(CurE->getMemberDecl())) { - if (!NoDiagnose) { - SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field) - << CurE->getSourceRange(); - return nullptr; - } - if (RelevantExpr) - return nullptr; - continue; + // 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()) { + if (!NoDiagnose) { + SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause) + << ME->getSourceRange() << getOpenMPClauseName(CKind); + return false; } + if (RelevantExpr) + return false; + return Visit(E->IgnoreParenImpCasts());; + } - auto *FD = cast(CurE->getMemberDecl()); + // 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().getNonReferenceType(); - // 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()) { - if (!NoDiagnose) { - SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause) - << CurE->getSourceRange() << getOpenMPClauseName(CKind); - return nullptr; - } - if (RelevantExpr) - return nullptr; - continue; + // 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 (CurType->isUnionType()) { + if (!NoDiagnose) { + SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed) + << ME->getSourceRange(); + return false; } + return Visit(E->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 = BaseE->getType().getNonReferenceType(); + // If we got a member expression, we should not expect any array section + // before that: + // + // 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. + // + AllowUnitySizeArraySection = false; + AllowWholeSizeArraySection = false; + + // Record the component. + Components.emplace_back(ME, FD); + if (RelevantExpr || Visit(E->IgnoreParenImpCasts())) { + Depth = CntStack.back(); + CntStack.pop_back(); + return true; + } else { + return false; + } + } + bool VisitCXXThisExpr(CXXThisExpr *CTE) { + if (!RelevantExpr) { + Components.emplace_back(CTE, nullptr); + this->RelevantExpr = cast(CTE); + } + return true; + } + bool VisitUnaryOperator(UnaryOperator *UO) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + if (UO->getOpcode() == UO_Deref || UO->getOpcode() == UO_AddrOf) { + if (!this->RelevantExpr && UO->getOpcode() == UO_Deref) { + Depth++; + // Record the component. + Components.emplace_back(UO, nullptr); + } else { + Depth--; + } + Expr *SubE = UO->getSubExpr()->IgnoreParenImpCasts(); + if (SubE && Visit(SubE->IgnoreParenImpCasts())) { + if (UO->getOpcode() == UO_Deref) + Depth--; + else + Depth++; + return true; + } + } + return false; + } + bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) { + Depth++; + Expr *E = AE->getBase()->IgnoreParenImpCasts(); - // 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 (CurType->isUnionType()) { - if (!NoDiagnose) { - SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed) - << CurE->getSourceRange(); - return nullptr; - } - continue; + if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) { + if (!NoDiagnose) { + SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) + << 0 << AE->getSourceRange(); + return false; } + return Visit(E->IgnoreParenImpCasts()); + } - // If we got a member expression, we should not expect any array section - // before that: - // - // 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. - // - AllowUnitySizeArraySection = false; + // If we got an array subscript that express the whole dimension we + // can have any array expressions before. If it only expressing part of + // the dimension, we can only have unitary-size array expressions. + if (checkArrayExpressionDoesNotReferToWholeSize(SemaRef, AE, + E->getType())) AllowWholeSizeArraySection = false; - // Record the component. - CurComponents.emplace_back(CurE, FD); - } else if (auto *CurE = dyn_cast(E)) { - E = CurE->getBase()->IgnoreParenImpCasts(); - - if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) { - if (!NoDiagnose) { - SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) - << 0 << CurE->getSourceRange(); - return nullptr; + if (const auto *TE = dyn_cast(E)) { + Expr::EvalResult Result; + if (AE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) { + if (!Result.Val.getInt().isNullValue()) { + SemaRef.Diag(AE->getIdx()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(AE->getIdx()->getExprLoc(), + diag::note_omp_invalid_subscript_on_this_ptr_map); } - continue; } + RelevantExpr = TE; + } - // If we got an array subscript that express the whole dimension we - // can have any array expressions before. If it only expressing part of - // the dimension, we can only have unitary-size array expressions. - if (checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, - E->getType())) - AllowWholeSizeArraySection = false; - - if (const auto *TE = dyn_cast(E)) { - Expr::EvalResult Result; - if (CurE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) { - if (!Result.Val.getInt().isNullValue()) { - SemaRef.Diag(CurE->getIdx()->getExprLoc(), - diag::err_omp_invalid_map_this_expr); - SemaRef.Diag(CurE->getIdx()->getExprLoc(), - diag::note_omp_invalid_subscript_on_this_ptr_map); - } - } - RelevantExpr = TE; - } + // Record the component - we don't have any declaration associated. + Components.emplace_back(AE, nullptr); - // Record the component - we don't have any declaration associated. - CurComponents.emplace_back(CurE, nullptr); - } else if (auto *CurE = dyn_cast(E)) { - assert(!NoDiagnose && "Array sections cannot be implicitly mapped."); - E = CurE->getBase()->IgnoreParenImpCasts(); + if (E && Visit(E->IgnoreParenImpCasts())) { + Depth--; + return true; + } + return false; + } + bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) { + assert(!NoDiagnose && "Array sections cannot be implicitly mapped."); + Depth++; + Expr *E = OASE->getBase()->IgnoreParenImpCasts(); - QualType CurType = - OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType(); + QualType CurType = + OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType(); - // 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 (CurType->isReferenceType()) - CurType = CurType->getPointeeType(); + // 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 (CurType->isReferenceType()) + CurType = CurType->getPointeeType(); - bool IsPointer = CurType->isAnyPointerType(); + bool IsPointer = CurType->isAnyPointerType(); - if (!IsPointer && !CurType->isArrayType()) { - SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) - << 0 << CurE->getSourceRange(); - return nullptr; - } + if (!IsPointer && !CurType->isArrayType()) { + SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name) + << 0 << OASE->getSourceRange(); + return false; + } - bool NotWhole = - checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, CurType); - bool NotUnity = - checkArrayExpressionDoesNotReferToUnitySize(SemaRef, CurE, CurType); + bool NotWhole = + checkArrayExpressionDoesNotReferToWholeSize(SemaRef, OASE, CurType); + bool NotUnity = + checkArrayExpressionDoesNotReferToUnitySize(SemaRef, OASE, CurType); - if (AllowWholeSizeArraySection) { - // Any array section is currently allowed. Allowing a whole size array - // section implies allowing a unity array section as well. - // - // If this array section refers to the whole dimension we can still - // accept other array sections before this one, except if the base is a - // pointer. Otherwise, only unitary sections are accepted. - if (NotWhole || IsPointer) - AllowWholeSizeArraySection = false; - } else if (AllowUnitySizeArraySection && NotUnity) { - // A unity or whole array section is not allowed and that is not - // compatible with the properties of the current array section. - SemaRef.Diag( - ELoc, diag::err_array_section_does_not_specify_contiguous_storage) - << CurE->getSourceRange(); - return nullptr; - } + if (AllowWholeSizeArraySection) { + // Any array section is currently allowed. Allowing a whole size array + // section implies allowing a unity array section as well. + // + // If this array section refers to the whole dimension we can still + // accept other array sections before this one, except if the base is a + // pointer. Otherwise, only unitary sections are accepted. + if (NotWhole || IsPointer) + AllowWholeSizeArraySection = false; + } else if (AllowUnitySizeArraySection && NotUnity) { + // A unity or whole array section is not allowed and that is not + // compatible with the properties of the current array section. + SemaRef.Diag( + ELoc, diag::err_array_section_does_not_specify_contiguous_storage) + << OASE->getSourceRange(); + return false; + } - if (const auto *TE = dyn_cast(E)) { - Expr::EvalResult ResultR; - Expr::EvalResult ResultL; - if (CurE->getLength()->EvaluateAsInt(ResultR, - SemaRef.getASTContext())) { - if (!ResultR.Val.getInt().isOneValue()) { - SemaRef.Diag(CurE->getLength()->getExprLoc(), - diag::err_omp_invalid_map_this_expr); - SemaRef.Diag(CurE->getLength()->getExprLoc(), - diag::note_omp_invalid_length_on_this_ptr_mapping); - } + if (const auto *TE = dyn_cast(E)) { + Expr::EvalResult ResultR; + Expr::EvalResult ResultL; + if (OASE->getLength()->EvaluateAsInt(ResultR, + SemaRef.getASTContext())) { + if (!ResultR.Val.getInt().isOneValue()) { + SemaRef.Diag(OASE->getLength()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(OASE->getLength()->getExprLoc(), + diag::note_omp_invalid_length_on_this_ptr_mapping); } - if (CurE->getLowerBound() && CurE->getLowerBound()->EvaluateAsInt( - ResultL, SemaRef.getASTContext())) { - if (!ResultL.Val.getInt().isNullValue()) { - SemaRef.Diag(CurE->getLowerBound()->getExprLoc(), - diag::err_omp_invalid_map_this_expr); - SemaRef.Diag(CurE->getLowerBound()->getExprLoc(), - diag::note_omp_invalid_lower_bound_on_this_ptr_mapping); - } + } + if (OASE->getLowerBound() && OASE->getLowerBound()->EvaluateAsInt( + ResultL, SemaRef.getASTContext())) { + if (!ResultL.Val.getInt().isNullValue()) { + SemaRef.Diag(OASE->getLowerBound()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(OASE->getLowerBound()->getExprLoc(), + diag::note_omp_invalid_lower_bound_on_this_ptr_mapping); } - RelevantExpr = TE; } + RelevantExpr = TE; + } - // Record the component - we don't have any declaration associated. - CurComponents.emplace_back(CurE, nullptr); - } else { - if (!NoDiagnose) { - // If nothing else worked, this is not a valid map clause expression. - SemaRef.Diag( - ELoc, diag::err_omp_expected_named_var_member_or_array_expression) - << ERange; - } - return nullptr; + // Record the component - we don't have any declaration associated. + Components.emplace_back(OASE, nullptr); + + if (RelevantExpr || Visit(E->IgnoreParenImpCasts())) { + Depth--; + return true; + } + return false; + } + bool VisitBinaryOperator(BinaryOperator *BO) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + if (BO->getOpcode() == BO_Assign) { + this->RelevantExpr = cast(BO); + // Record the component. + Components.emplace_back(BO, nullptr); + return true; } + Expr *LE = BO->getLHS()->IgnoreParenImpCasts(); + Expr *RE = BO->getRHS()->IgnoreParenImpCasts(); + if (!LE && !RE) + return false; + if (LE && RE) { + Visit(LE); + Visit(RE); + return true; + } + return false; + } + bool VisitConditionalOperator(ConditionalOperator *CO) { + emitErrorMsg(); + return false; } + bool VisitCallExpr(CallExpr *CE) { + emitErrorMsg(); + return false; + } + bool VisitCastExpr(CastExpr *CE) { + if (SemaRef.getLangOpts().OpenMP < 50) { + emitErrorMsg(); + return false; + } + if (!CE->isLValue()) + return false; + return Visit(CE->getSubExpr()); + } + bool VisitCXXMemberCallExpr(CXXMemberCallExpr *CCE) { + emitErrorMsg(); + return false; + } + const Expr *getFoundBase() const { + return RelevantExpr; + } + OMPClauseMappableExprCommon::MappableExprComponentList getComponents() { + return Components; + } + explicit MapBaseChecker( + Sema &SemaRef, OpenMPClauseKind CKind, + OMPClauseMappableExprCommon::MappableExprComponentList &Components, + bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange) + : SemaRef(SemaRef), CKind(CKind), Components(Components), + NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {} +}; +} // namespace - return RelevantExpr; +// Return the expression of the base of the mappable expression or null if it +// cannot be determined and do all the necessary checks to see if the expression +// is valid as a standalone mappable expression. In the process, record all the +// components of the expression. +static const Expr *checkMapClauseExpressionBase( + Sema &SemaRef, Expr *E, + OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, + OpenMPClauseKind CKind, bool NoDiagnose) { + SourceLocation ELoc = E->getExprLoc(); + SourceRange ERange = E->getSourceRange(); + MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc, + ERange); + if (Checker.Visit(E->IgnoreParenImpCasts())) { + return Checker.getFoundBase(); + } else { + return nullptr; + } } // Return true if expression E associated with value VD has conflicts with other @@ -15898,11 +16053,19 @@ } Expr *SimpleExpr = RE->IgnoreParenCasts(); + Expr::isModifiableLvalueResult IsLV = + VE->isModifiableLvalue(SemaRef.Context, &ELoc); - if (!RE->IgnoreParenImpCasts()->isLValue()) { - SemaRef.Diag(ELoc, - diag::err_omp_expected_named_var_member_or_array_expression) - << RE->getSourceRange(); + if (!RE->IgnoreParenImpCasts()->isLValue() && + IsLV != Expr::MLV_LValueCast) { + if (SemaRef.getLangOpts().OpenMP < 50) { + SemaRef.Diag( + ELoc, diag::err_omp_expected_named_var_member_or_array_expression) + << RE->getSourceRange(); + } else { + SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses) + << RE->getSourceRange(); + } continue; } @@ -15937,6 +16100,15 @@ MVLI.VarBaseDeclarations.push_back(nullptr); continue; } + if (const auto *BO = dyn_cast(BE)) { + MVLI.UDMapperList.push_back(nullptr); + MVLI.ProcessedVarList.push_back(RE); + MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); + MVLI.VarComponents.back().append(CurComponents.begin(), + CurComponents.end()); + MVLI.VarBaseDeclarations.push_back(nullptr); + continue; + } // For the following checks, we rely on the base declaration which is // expected to be associated with the last component. The declaration is diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp --- a/clang/test/OpenMP/target_map_messages.cpp +++ b/clang/test/OpenMP/target_map_messages.cpp @@ -1,7 +1,13 @@ -// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -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 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -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 // RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized #ifdef CCODE void foo(int arg) { @@ -57,9 +63,9 @@ {} #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,a*2) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} {} - #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,(c+1)[2]) // le45-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}} {} @@ -305,7 +311,7 @@ {} #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.S.foo()[:12]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} {} #pragma omp target map(r.C, r.D) {} @@ -339,7 +345,7 @@ {} #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((p+1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} #pragma omp target map(u.B) // expected-error {{mapping of union members is not allowed}} {} @@ -479,7 +485,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 expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target map(I) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} foo(); #pragma omp target map(S2::S2s) foo(); @@ -496,7 +502,7 @@ #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) // expected-error 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) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} #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}} @@ -592,7 +598,7 @@ #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]) // expected-error {{xpected 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]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} #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}} diff --git a/clang/test/OpenMP/target_teams_map_messages.cpp b/clang/test/OpenMP/target_teams_map_messages.cpp --- a/clang/test/OpenMP/target_teams_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_map_messages.cpp @@ -1,10 +1,14 @@ // RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -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 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized // RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %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 +// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized #ifdef CCODE void foo(int arg) { const int n = 0; @@ -44,9 +48,9 @@ {} #pragma omp target teams map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}} {} - #pragma omp target teams map(arg,a*2) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + #pragma omp target teams map(arg,a*2) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} - #pragma omp target teams map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + #pragma omp target teams map(arg,(c+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} #pragma omp target teams map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}} {} @@ -252,7 +256,7 @@ {} #pragma omp target teams map(r.S.Arr[:12]) {} - #pragma omp target teams map(r.S.foo()[:12]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + #pragma omp target teams map(r.S.foo()[:12]) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} #pragma omp target teams map(r.C, r.D) {} @@ -286,7 +290,7 @@ {} #pragma omp target teams map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target teams map((p+1)->A) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} + #pragma omp target teams map((p+1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} {} #pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}} {} @@ -416,7 +420,7 @@ foo(); #pragma omp target teams map(T) // expected-error {{'T' does not refer to a value}} foo(); -#pragma omp target teams map(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target teams map(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} foo(); #pragma omp target teams map(S2::S2s) foo(); @@ -434,7 +438,8 @@ 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) // expected-error 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) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-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}} @@ -508,7 +513,7 @@ 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]) // expected-error {{xpected 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]) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-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}} diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -291,5 +291,405 @@ #pragma omp target update from(arg) if(arg) device(4) {++arg;} } +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 + +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +#ifdef CK5 + +// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33] + +// CK5-LABEL: lvalue +void lvalue(int *B, int l, int e) { + + // CK5-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** + // CK5-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]] + // CK5-DAG: store i32* [[ONE:%.+]], i32** [[PC0]] + // CK5-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]] + // CK5-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]] + #pragma omp target update to(*B) + *B += e; + #pragma omp target update from(*B) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-64 +// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-64 +// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-32 +// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-32 + +// RUN: %clang_cc1 -DCK6 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK6 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK6 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK6 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +#ifdef CK6 + +// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33] + +// CK6-LABEL: lvalue +void lvalue(int *B, int l, int e) { + + // CK6-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** + // CK6-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]] + // CK6-DAG: store i32* [[ADD_PTR:%.+]], i32** [[PC0]] + // CK6-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]] + // CK6-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]] + // CK6-DAG: store i32 [[L:%.+]], i32* [[L_ADDR]] + #pragma omp target update to(*(B+l)) + *(B+l) += e; + #pragma omp target update from(*(B+l)) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-64 +// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-64 +// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-32 +// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-32 + +// RUN: %clang_cc1 -DCK7 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK7 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK7 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK7 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +#ifdef CK7 + +// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33] + +// CK7-LABEL: lvalue +void lvalue(int *B, int l, int e) { + + // CK7-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** + // CK7-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]] + // CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]] + // CK7-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]] + // CK7-DAG: [[ARRAY_IDX]] = getelementptr inbounds i32, i32* [[ADD_PTR:%.+]], i{{32|64}} [[IDX_PROM:%.+]] + // CK7-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i64 [[IDX_EXT:%.+]] + // CK7-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i32 [[TWO:%.+]] + // CK7-32-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]] + // CK7-32-DAG: [[TWO]] = load i32, i32* [[L_ADDR:%.+]] + // CK7-32-DAG: [[IDX_PROM]] = load i32, i32* [[L_ADDR]] + // CK7-64-DAG: [[IDX_EXT:%.+]] = sext i32 [[TWO:%.+]] to i64 + // CK7-64-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]] + // CK7-64-DAG: [[IDX_PROM]] = sext i32 [[THREE:%.+]] to i64 + // CK7-64-DAG: [[THREE]] = load i32, i32* [[L_ADDR]] + #pragma omp target update to((B+l)[l]) + (B+l)[l] += e; + #pragma omp target update from((B+l)[l]) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-64 +// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-64 +// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-32 +// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-32 + +// RUN: %clang_cc1 -DCK8 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK8 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK8 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK8 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +#ifdef CK8 +// CK8-64: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 8, i64 4] +// CK8-32: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 4, i64 4] +// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] + +// CK8-LABEL: lvalue +void lvalue(int **B, int l, int e) { + + // CK8-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}], [2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK8-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK8-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK8-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32*** + // CK8-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** + // CK8-DAG: store i32** [[ARRAY_IDX_1:%.+]], i32*** [[BPC0]] + // CK8-DAG: store i32* [[ARRAY_IDX_4:%.+]], i32** [[PC0]] + // CK8-DAG: store i32** [[ARRAY_IDX_1]], i32*** [[NINE:%.+]] + // CK8-DAG: [[ARRAY_IDX_1]] = getelementptr inbounds i32*, i32** [[ADD_PTR:%.+]], i{{.+}} 1 + // CK8-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[IDX_EXT:%.+]] + // CK8-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64 + // CK8-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2 + // CK8-64-DAG: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1 + // CK8-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[IDX_EXT_1:%.+]] + // CK8-64-DAG: [[IDX_EXT_1]] = sext i32 [[L_VAL:%.+]] to i{{.+}} + // CK8-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[L_VAL:%.+]] + // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1 + // CK8-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[L_VAL:%.+]] + #pragma omp target update to((B+l)[1][2]) + (B+l)[1][2] += e; + #pragma omp target update from((B+l)[1][2]) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-64 +// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-64 +// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-32 +// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-32 + +// RUN: %clang_cc1 -DCK9 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK9 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK9 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK9 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} + +#ifdef CK9 + +struct S { + double *p; +}; + +// CK9: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 32, i64 281474976710673] +// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674] + +// CK9-LABEL: lvalue +void lvalue(struct S *s, int l, int e) { + + // CK9-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] + // + // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 + // CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** + // CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** + // CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]] + // CK9-DAG: store double* [[THREE:%.+]], double** [[PC0]] + // CK9-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]] + // CK9-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 0 + // CK9-DAG: [[ONE]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]] + // CK9-DAG: [[THREE]] = load double*, double** [[P_1:%.+]], + // CK9-DAG: [[P_1]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[TWO:%.+]], i32 0, i32 0 + // CK9-DAG: [[TWO]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]] + #pragma omp target update to(*(s->p)) + *(s->p) += e; + #pragma omp target update from(*(s->p)) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-64 +// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-64 +// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-32 +// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-32 + +// RUN: %clang_cc1 -DCK10 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK10 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK10 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK10 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +#ifdef CK10 + +struct S { + double *p; +}; + +// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674] + +// CK10-LABEL: lvalue +void lvalue(struct S *s, int l, int e) { + + // CK10-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] + // + // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 + // CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** + // CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** + // CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]] + // CK10-DAG: store double* [[ADD_PTR:%.+]], double** [[PC0]] + // CK10-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]] + // CK10-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]] + // CK10-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]] + // CK10-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64 + // CK10-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]] + // CK10-DAG: getelementptr inbounds {{.+}}, {{.+}}* [[SS_1:%.+]], i32 0, i32 0 + // CK10-DAG: [[SS_1]] = load {{.+}}*, {{.+}}** [[S_ADDR:%.+]] + #pragma omp target update to(*(s->p+l)) + *(s->p+l) += e; + #pragma omp target update from(*(s->p+l)) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-64 +// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-64 +// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-32 +// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-32 + +// RUN: %clang_cc1 -DCK11 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK11 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK11 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK11 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +#ifdef CK11 + +struct S { + double *p; +}; +// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674] + +// CK11-LABEL: lvalue +void lvalue(struct S *s, int l, int e) { + + // CK11-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] + // + // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 + // CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** + // CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** + // CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]] + // CK11-DAG: store double* [[ARRAY_IDX:%.+]], double** [[PC0]] + // CK11-DAG: store i{{.+}} 8, i{{.+}} [[SIZE0]] + // CK11-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[SS_1:%.+]], i32 0, i32 0 + // CK11-DAG: [[ARRAY_IDX]] = getelementptr inbounds double, double* [[ADD_PTR:%.+]], i{{.+}} 3 + // CK11-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]] + // CK11-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]] + // CK11-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64 + #pragma omp target update to((s->p+l)[3]) + (s->p+l)[3] += e; + #pragma omp target update from((s->p+l)[3]) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-64 +// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-64 +// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-32 +// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-32 + +// RUN: %clang_cc1 -DCK12 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK12 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK12 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK12 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +#ifdef CK12 + +struct S { + double *p; + struct S *sp; +}; +// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17] + +// CK12-LABEL: lvalue +void lvalue(struct S *s, int l, int e) { + + // CK12-DAG: call void @__tgt_target_data_update(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] + // + // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double*** + // CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double** + // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]] + // CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]] + // CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]] + // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]*** + // CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double*** + // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]] + // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]] + // CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]] + // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0 + // CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]** + // CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]*** + // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]] + // CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]] + // CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1 + #pragma omp target update to(*(s->sp->p)) + *(s->sp->p) = e; + #pragma omp target update from(*(s->sp->p)) +} + #endif #endif diff --git a/clang/test/OpenMP/target_update_from_messages.cpp b/clang/test/OpenMP/target_update_from_messages.cpp --- a/clang/test/OpenMP/target_update_from_messages.cpp +++ b/clang/test/OpenMP/target_update_from_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -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 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized void foo() { } @@ -59,6 +62,30 @@ double *p; unsigned bfa : 4; }; +struct S8 { + int *ptr; + int a; + struct S7* S; + void foo() { +#pragma omp target update from(*this) // 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))) // 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)) // 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->S->i+this->S->p)))) // 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->S->i+this->S->s6[0].pp)))) // 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(*(&(*(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'}} + {} + } +}; + +struct Base { + virtual ~Base() {} +}; +struct Derived: Base { + virtual void name() {} +}; + S3 h; #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}} @@ -89,12 +116,13 @@ #pragma omp target update from(x) #pragma omp target update from(t[:I]) #pragma omp target update from(T) // expected-error {{'T' 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(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update from(S2::S2s) #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) // expected-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) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} 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}} #pragma omp target update from(ba) @@ -106,7 +134,7 @@ #pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} #pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} -#pragma omp target update from(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(x, (m+1)[2]) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update from(s7.i, s7.a[:3]) #pragma omp target update from(s7.s6[1].aa[0:5]) #pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} @@ -119,6 +147,16 @@ #pragma omp target update from(s7.x) } + const int &rci = i; + Base *b = new Base; + int *pd = const_cast(&d); +#pragma omp target update to(const_cast(rci)) +#pragma omp target update to(const_cast(pd)) +#pragma omp target update to(static_cast(i)) +#pragma omp target update to(dynamic_cast(b)) +#pragma omp target update to(reinterpret_cast(b)) + {} + return 0; } @@ -136,6 +174,8 @@ const int (&l)[5] = da; int *m; S7 s7; + int **BB, *offset; + int ***BBB; #pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} @@ -148,7 +188,7 @@ #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) // expected-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) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} 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(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) @@ -160,7 +200,7 @@ #pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} #pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} -#pragma omp target update from(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(x, (m+1)[2]) // // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update from(s7.i, s7.a[:3]) #pragma omp target update from(s7.s6[1].aa[0:5]) #pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} @@ -168,11 +208,35 @@ #pragma omp target update from(s7.p[:10]) #pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}} #pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target update from(*(&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(*(&(*(&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(**(BB+*offset)) // 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(**(BB+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(*(m+*offset)) // 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(**(*offset+BB)) // 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(**(y+BB)) // 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(*(*offset+m)) // 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(**(*offset+BB+*m)) // 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(**(*(*(&offset))+BB+*m)) // 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+*(**BB+BBB)+s7.i)+static_cast(*(s7.s6[0].pp)))) // 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 data map(to: s7.i) { #pragma omp target update from(s7.x) } + const int &rci = i; + Base *b = new Base; + int *pd = const_cast(&d); +#pragma omp target update from(const_cast(rci)) +#pragma omp target update from(const_cast(pd)) +#pragma omp target update from(static_cast(i)) +#pragma omp target update from(dynamic_cast(b)) +#pragma omp target update from(reinterpret_cast(b)) + {} + +#pragma omp target update from(static_cast(i + 1)) // 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'}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} + {} + 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}} } diff --git a/clang/test/OpenMP/target_update_to_messages.cpp b/clang/test/OpenMP/target_update_to_messages.cpp --- a/clang/test/OpenMP/target_update_to_messages.cpp +++ b/clang/test/OpenMP/target_update_to_messages.cpp @@ -1,6 +1,12 @@ -// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -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 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -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 void foo() { } @@ -59,6 +65,29 @@ double *p; unsigned bfa : 4; }; +struct S8 { + int *ptr; + int a; + struct S7* S; + void foo() { +#pragma omp target update to(*this) // 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(*(&(*this))) // 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(*(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 to(*(&(*(this->S->i+this->S->p)))) // 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(*(&(*(this->S->i+this->S->s6[0].pp)))) // 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(*(&(*(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 to(*(&(*(*(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'}} + {} + } +}; + +struct Base { + virtual ~Base() {} +}; +struct Derived: Base { + virtual void name() {} +}; S3 h; #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}} @@ -89,12 +118,12 @@ #pragma omp target update to(x) #pragma omp target update to(t[:I]) #pragma omp target update to(T) // expected-error {{'T' 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(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update to(S2::S2s) #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) // expected-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) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} 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) @@ -106,7 +135,7 @@ #pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} #pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} -#pragma omp target update to(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(x, (m+1)[2]) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update to(s7.i, s7.a[:3]) #pragma omp target update to(s7.s6[1].aa[0:5]) #pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} @@ -118,6 +147,16 @@ { #pragma omp target update to(s7.x) } + + const int &rci = i; + Base *b = new Base; + int *pd = const_cast(&d); +#pragma omp target update to(const_cast(rci)) +#pragma omp target update to(const_cast(pd)) +#pragma omp target update to(static_cast(i)) +#pragma omp target update to(dynamic_cast(b)) +#pragma omp target update to(reinterpret_cast(b)) + {} return 0; } @@ -135,6 +174,8 @@ const int (&l)[5] = da; S7 s7; int *m; + int **BB, *offset; + int ***BBB; #pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} @@ -147,7 +188,7 @@ #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) // expected-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) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} 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(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) @@ -159,7 +200,7 @@ #pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} #pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} -#pragma omp target update to(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(x, (m+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} #pragma omp target update to(s7.i, s7.a[:3]) #pragma omp target update to(s7.s6[1].aa[0:5]) #pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} @@ -172,6 +213,32 @@ #pragma omp target update to(s7.x) } +#pragma omp target update to(*(&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(*(&(*(&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(**(BB+*offset)) // 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(**(BB+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(*(m+*offset)) // 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(**(*offset+BB)) // 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(**(y+BB)) // 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(*(*offset+m)) // 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(**(*offset+BB+*m)) // 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(**(*(*(&offset))+BB+*m)) // 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+*(**BB+BBB)+s7.i)+static_cast(*(s7.s6[0].pp)))) // 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'}} + {} + + const int &rci = i; + Base *b = new Base; + int *pd = const_cast(&d); +#pragma omp target update to(const_cast(rci)) +#pragma omp target update to(const_cast(pd)) +#pragma omp target update to(static_cast(i)) +#pragma omp target update to(dynamic_cast(b)) +#pragma omp target update to(reinterpret_cast(b)) + {} + +#pragma omp target update to(static_cast(i + 1)) // 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'}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} + {} + 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}} }