Index: clang/include/clang/Sema/ScopeInfo.h =================================================================== --- clang/include/clang/Sema/ScopeInfo.h +++ clang/include/clang/Sema/ScopeInfo.h @@ -756,13 +756,16 @@ unsigned short CapRegionKind; unsigned short OpenMPLevel; + unsigned short OpenMPCaptureLevel; CapturedRegionScopeInfo(DiagnosticsEngine &Diag, Scope *S, CapturedDecl *CD, RecordDecl *RD, ImplicitParamDecl *Context, - CapturedRegionKind K, unsigned OpenMPLevel) + CapturedRegionKind K, unsigned OpenMPLevel, + unsigned OpenMPCaptureLevel) : CapturingScopeInfo(Diag, ImpCap_CapturedRegion), TheCapturedDecl(CD), TheRecordDecl(RD), TheScope(S), - ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel) { + ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel), + OpenMPCaptureLevel(OpenMPCaptureLevel) { Kind = SK_CapturedRegion; } Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -1415,8 +1415,8 @@ void RecordParsingTemplateParameterDepth(unsigned Depth); void PushCapturedRegionScope(Scope *RegionScope, CapturedDecl *CD, - RecordDecl *RD, - CapturedRegionKind K); + RecordDecl *RD, CapturedRegionKind K, + unsigned OpenMPCaptureLevel = 0); /// Custom deleter to allow FunctionScopeInfos to be kept alive for a short /// time after they've been popped. @@ -3973,7 +3973,8 @@ typedef std::pair CapturedParamNameType; void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope, CapturedRegionKind Kind, - ArrayRef Params); + ArrayRef Params, + unsigned OpenMPCaptureLevel = 0); StmtResult ActOnCapturedRegionEnd(Stmt *S); void ActOnCapturedRegionError(); RecordDecl *CreateCapturedStmtRecordDecl(CapturedDecl *&CD, @@ -9021,7 +9022,9 @@ /// reference. /// \param Level Relative level of nested OpenMP construct for that the check /// is performed. - bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const; + /// \param OpenMPCaptureLevel Capture level within an OpenMP construct. + bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, + unsigned OpenMPCaptureLevel) const; /// Check if the specified variable is used in one of the private /// clauses (private, firstprivate, lastprivate, reduction etc.) in OpenMP Index: clang/lib/Sema/Sema.cpp =================================================================== --- clang/lib/Sema/Sema.cpp +++ clang/lib/Sema/Sema.cpp @@ -2105,10 +2105,12 @@ } void Sema::PushCapturedRegionScope(Scope *S, CapturedDecl *CD, RecordDecl *RD, - CapturedRegionKind K) { - CapturingScopeInfo *CSI = new CapturedRegionScopeInfo( + CapturedRegionKind K, + unsigned OpenMPCaptureLevel) { + auto *CSI = new CapturedRegionScopeInfo( getDiagnostics(), S, CD, RD, CD->getContextParam(), K, - (getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0); + (getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0, + OpenMPCaptureLevel); CSI->ReturnType = Context.VoidTy; FunctionScopes.push_back(CSI); } Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -15660,7 +15660,8 @@ if (HasConst) DeclRefType.addConst(); } - ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel); + ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel, + RSI->OpenMPCaptureLevel); } if (ByRef) Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -522,6 +522,13 @@ assert(!isStackEmpty() && "No directive at specified level."); return getStackElemAtLevel(Level).Directive; } + /// Returns the capture region at the specified level. + OpenMPDirectiveKind getCaptureRegion(unsigned Level, + unsigned OpenMPCaptureLevel) const { + SmallVector CaptureRegions; + getOpenMPCaptureRegions(CaptureRegions, getDirective(Level)); + return CaptureRegions[OpenMPCaptureLevel]; + } /// Returns parent directive. OpenMPDirectiveKind getParentDirective() const { const SharingMapTy *Parent = getSecondOnStackOrNull(); @@ -1340,7 +1347,8 @@ } const_iterator End = end(); if (!SemaRef.isOpenMPCapturedByRef( - D, std::distance(ParentIterTarget, End))) { + D, std::distance(ParentIterTarget, End), + /*OpenMPCaptureLevel=*/0)) { DVar.RefExpr = buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(), IterTarget->ConstructLoc); @@ -1606,7 +1614,8 @@ << Context.getTargetInfo().getTriple().str() << E->getSourceRange(); } -bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const { +bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level, + unsigned OpenMPCaptureLevel) const { assert(LangOpts.OpenMP && "OpenMP is not allowed"); ASTContext &Ctx = getASTContext(); @@ -1616,6 +1625,7 @@ D = cast(D->getCanonicalDecl()); QualType Ty = D->getType(); + bool IsVariableUsedInMapClause = false; if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective, Level)) { // This table summarizes how a given variable should be passed to the device // given its type and the clauses where it appears. This table is based on @@ -1677,7 +1687,6 @@ // Locate map clauses and see if the variable being captured is referred to // in any of those clauses. Here we only care about variables, not fields, // because fields are part of aggregates. - bool IsVariableUsedInMapClause = false; bool IsVariableAssociatedWithSection = false; DSAStack->checkMappableExprComponentListsForDeclAtLevel( @@ -1735,10 +1744,13 @@ if (IsByRef && Ty.getNonReferenceType()->isScalarType()) { IsByRef = - !DSAStack->hasExplicitDSA( - D, - [](OpenMPClauseKind K) -> bool { return K == OMPC_firstprivate; }, - Level, /*NotLastprivate=*/true) && + ((IsVariableUsedInMapClause && + DSAStack->getCaptureRegion(Level, OpenMPCaptureLevel) == + OMPD_target) || + !DSAStack->hasExplicitDSA( + D, + [](OpenMPClauseKind K) -> bool { return K == OMPC_firstprivate; }, + Level, /*NotLastprivate=*/true)) && // If the variable is artificial and must be captured by value - try to // capture by value. !(isa(D) && !D->hasAttr() && @@ -2960,7 +2972,7 @@ std::make_pair(StringRef(), QualType()) // __context with shared vars }; ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - Params); + Params, /*OpenMPCaptureLevel=*/0); // Mark this captured region as inlined, because we don't use outlined // function directly. getCurCapturedRegion()->TheCapturedDecl->addAttr( @@ -2971,7 +2983,7 @@ }; // Start a captured region for 'target' with no implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsTarget); + ParamsTarget, /*OpenMPCaptureLevel=*/1); Sema::CapturedParamNameType ParamsTeamsOrParallel[] = { std::make_pair(".global_tid.", KmpInt32PtrTy), std::make_pair(".bound_tid.", KmpInt32PtrTy), @@ -2980,7 +2992,7 @@ // Start a captured region for 'teams' or 'parallel'. Both regions have // the same implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsTeamsOrParallel); + ParamsTeamsOrParallel, /*OpenMPCaptureLevel=*/2); break; } case OMPD_target: @@ -3004,14 +3016,15 @@ std::make_pair(StringRef(), QualType()) // __context with shared vars }; ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - Params); + Params, /*OpenMPCaptureLevel=*/0); // Mark this captured region as inlined, because we don't use outlined // function directly. getCurCapturedRegion()->TheCapturedDecl->addAttr( AlwaysInlineAttr::CreateImplicit( Context, AlwaysInlineAttr::Keyword_forceinline)); ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - std::make_pair(StringRef(), QualType())); + std::make_pair(StringRef(), QualType()), + /*OpenMPCaptureLevel=*/1); break; } case OMPD_simd: @@ -3143,7 +3156,7 @@ std::make_pair(StringRef(), QualType()) // __context with shared vars }; ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - Params); + Params, /*OpenMPCaptureLevel=*/0); // Mark this captured region as inlined, because we don't use outlined // function directly. getCurCapturedRegion()->TheCapturedDecl->addAttr( @@ -3154,7 +3167,7 @@ }; // Start a captured region for 'target' with no implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsTarget); + ParamsTarget, /*OpenMPCaptureLevel=*/1); Sema::CapturedParamNameType ParamsTeams[] = { std::make_pair(".global_tid.", KmpInt32PtrTy), @@ -3163,7 +3176,7 @@ }; // Start a captured region for 'target' with no implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsTeams); + ParamsTeams, /*OpenMPCaptureLevel=*/2); Sema::CapturedParamNameType ParamsParallel[] = { std::make_pair(".global_tid.", KmpInt32PtrTy), @@ -3175,7 +3188,7 @@ // Start a captured region for 'teams' or 'parallel'. Both regions have // the same implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsParallel); + ParamsParallel, /*OpenMPCaptureLevel=*/3); break; } @@ -3192,7 +3205,7 @@ }; // Start a captured region for 'target' with no implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsTeams); + ParamsTeams, /*OpenMPCaptureLevel=*/0); Sema::CapturedParamNameType ParamsParallel[] = { std::make_pair(".global_tid.", KmpInt32PtrTy), @@ -3204,7 +3217,7 @@ // Start a captured region for 'teams' or 'parallel'. Both regions have // the same implicit parameters. ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP, - ParamsParallel); + ParamsParallel, /*OpenMPCaptureLevel=*/1); break; } case OMPD_target_update: @@ -11252,7 +11265,13 @@ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] // A list item cannot appear in both a map clause and a data-sharing // attribute clause on the same construct - if (isOpenMPTargetExecutionDirective(CurrDir)) { + // + // OpenMP 5.0 [2.19.7.1, Restrictions, p.7] + // A list item cannot appear in both a map clause and a data-sharing + // attribute clause on the same construct unless the construct is a + // combined construct. + if ((LangOpts.OpenMP <= 45 && isOpenMPTargetExecutionDirective(CurrDir)) || + CurrDir == OMPD_target) { OpenMPClauseKind ConflictKind; if (DSAStack->checkMappableExprComponentListsForDecl( VD, /*CurrentRegionOnly=*/true, @@ -11487,7 +11506,14 @@ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] // A list item cannot appear in both a map clause and a data-sharing // attribute clause on the same construct - if (isOpenMPTargetExecutionDirective(CurrDir)) { + // + // OpenMP 5.0 [2.19.7.1, Restrictions, p.7] + // A list item cannot appear in both a map clause and a data-sharing + // attribute clause on the same construct unless the construct is a + // combined construct. + if ((LangOpts.OpenMP <= 45 && + isOpenMPTargetExecutionDirective(CurrDir)) || + CurrDir == OMPD_target) { OpenMPClauseKind ConflictKind; if (DSAStack->checkMappableExprComponentListsForDecl( VD, /*CurrentRegionOnly=*/true, @@ -14601,7 +14627,14 @@ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] // A list item cannot appear in both a map clause and a data-sharing // attribute clause on the same construct - if (VD && isOpenMPTargetExecutionDirective(DKind)) { + // + // OpenMP 5.0 [2.19.7.1, Restrictions, p.7] + // A list item cannot appear in both a map clause and a data-sharing + // attribute clause on the same construct unless the construct is a + // combined construct. + if (VD && ((SemaRef.LangOpts.OpenMP <= 45 && + isOpenMPTargetExecutionDirective(DKind)) || + DKind == OMPD_target)) { DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false); if (isOpenMPPrivate(DVar.CKind)) { SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa) Index: clang/lib/Sema/SemaStmt.cpp =================================================================== --- clang/lib/Sema/SemaStmt.cpp +++ clang/lib/Sema/SemaStmt.cpp @@ -4334,7 +4334,8 @@ void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope, CapturedRegionKind Kind, - ArrayRef Params) { + ArrayRef Params, + unsigned OpenMPCaptureLevel) { CapturedDecl *CD = nullptr; RecordDecl *RD = CreateCapturedStmtRecordDecl(CD, Loc, Params.size()); @@ -4379,7 +4380,7 @@ CD->setContextParam(ParamNum, Param); } // Enter the capturing scope for this captured region. - PushCapturedRegionScope(CurScope, CD, RD, Kind); + PushCapturedRegionScope(CurScope, CD, RD, Kind, OpenMPCaptureLevel); if (CurScope) PushDeclContext(CurScope, CD); Index: clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized typedef void **omp_allocator_handle_t; extern const omp_allocator_handle_t omp_default_mem_alloc; @@ -240,7 +243,7 @@ #pragma omp target teams distribute parallel for lastprivate(si) // OK for (i = 0; i < argc; ++i) si = i + 1; -#pragma omp target teams distribute parallel for lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} expected-note {{defined as lastprivate}} +#pragma omp target teams distribute parallel for lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} le45-note {{defined as lastprivate}} for (i = 0; i < argc; ++i) foo(); return foomain(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain' requested here}} Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized typedef void **omp_allocator_handle_t; extern const omp_allocator_handle_t omp_default_mem_alloc; @@ -240,7 +243,7 @@ #pragma omp target teams distribute parallel for simd lastprivate(si) // OK for (i = 0; i < argc; ++i) si = i + 1; -#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note {{defined as lastprivate}} +#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} le45-note {{defined as lastprivate}} for (i = 0; i < argc; ++i) foo(); return foomain(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain' requested here}} } Index: clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized typedef void **omp_allocator_handle_t; extern const omp_allocator_handle_t omp_default_mem_alloc; @@ -131,8 +134,10 @@ #pragma omp target teams distribute simd firstprivate(h) // expected-error {{threadprivate or thread local variable cannot be firstprivate}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note 2 {{defined as private}} - for (i = 0; i < argc; ++i) foo(); // expected-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}} +#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note {{defined as private}} + // le45-note@-1 {{defined as private}} + // FIXME: Error is dropped for OpenMP 5.0. Probably shouldn't be. + for (i = 0; i < argc; ++i) foo(); // le45-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}} #pragma omp target teams distribute simd firstprivate(i) for (j = 0; j < argc; ++j) foo(); @@ -147,7 +152,7 @@ #pragma omp target teams distribute simd lastprivate(argc), firstprivate(argc) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd firstprivate(argc) map(argc) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as firstprivate}} +#pragma omp target teams distribute simd firstprivate(argc) map(argc) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as firstprivate}} for (i = 0; i < argc; ++i) foo(); return 0; Index: clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized typedef void **omp_allocator_handle_t; extern const omp_allocator_handle_t omp_default_mem_alloc; @@ -240,7 +243,7 @@ #pragma omp target teams distribute simd lastprivate(si) // OK for (i = 0; i < argc; ++i) si = i + 1; -#pragma omp target teams distribute simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as lastprivate}} +#pragma omp target teams distribute simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as lastprivate}} for (i = 0; i < argc; ++i) foo(); return foomain(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain' requested here}} Index: clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp @@ -1,6 +1,9 @@ -// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized typedef void **omp_allocator_handle_t; extern const omp_allocator_handle_t omp_default_mem_alloc; @@ -132,7 +135,7 @@ for (int k = 0; k < 10; ++k) { } -#pragma omp target teams distribute simd private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as private}} +#pragma omp target teams distribute simd private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as private}} for (int k = 0; k < argc; ++k) ++k; return 0; Index: clang/test/OpenMP/target_teams_map_codegen.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_teams_map_codegen.cpp @@ -0,0 +1,172 @@ +// Test host codegen. +// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128 +// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128 +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST + +// Test target codegen - host bc file has to be created first. +// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes=CHECK,INT128 +// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s +// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,INT128 +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s + +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +// HOST: @[[MAPTYPES_FIRSTPRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] +// HOST: @[[MAPTYPES_REDUCTION:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] +// HOST: @[[MAPTYPES_FROM:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 34] +// HOST: @[[MAPTYPES_TO:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 33] +// HOST: @[[MAPTYPES_ALLOC:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 32] +// HOST: @[[MAPTYPES_ARRAY_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] +// HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 33, i64 33] +// HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] +// HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 34, i64 34] +// +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_PRIVATE:__omp_offloading_[^"\\]*mapWithPrivate[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FIRSTPRIVATE:__omp_offloading_[^"\\]*mapWithFirstprivate[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_REDUCTION:__omp_offloading_[^"\\]*mapWithReduction[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FROM:__omp_offloading_[^"\\]*mapFrom[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_TO:__omp_offloading_[^"\\]*mapTo[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ALLOC:__omp_offloading_[^"\\]*mapAlloc[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R0:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00" +// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R1:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00" +// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R0:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00" +// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R1:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00" + +// HOST: define {{.*}}mapWithPrivate +// HOST: call {{.*}} @.[[OFFLOAD_PRIVATE]].region_id +// HOST-NOT: offload_maptypes +// HOST-SAME: {{$}} +// +// CHECK: define {{.*}} void @[[OFFLOAD_PRIVATE]]() +// CHECK: call void ({{.*}}@[[OUTLINE_PRIVATE:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_PRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid.) +void mapWithPrivate() { + int x, y; + #pragma omp target teams private(x) map(x,y) private(y) + ; +} + +// HOST: define {{.*}}mapWithFirstprivate +// HOST: call {{.*}} @.[[OFFLOAD_FIRSTPRIVATE]].region_id{{.*}} @[[MAPTYPES_FIRSTPRIVATE]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) +// CHECK: call void ({{.*}}@[[OUTLINE_FIRSTPRIVATE:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_FIRSTPRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x, i{{[0-9]*}} %y) +void mapWithFirstprivate() { + int x, y; + #pragma omp target teams firstprivate(x) map(x,y) firstprivate(y) + ; +} + +// HOST: define {{.*}}mapWithReduction +// HOST: call {{.*}} @.[[OFFLOAD_REDUCTION]].region_id{{.*}} @[[MAPTYPES_REDUCTION]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) +// CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) +// CHECK: %.omp.reduction.red_list = alloca [2 x i8*] +void mapWithReduction() { + int x, y; + #pragma omp target teams reduction(+:x) map(x,y) reduction(+:y) + ; +} + +// HOST: define {{.*}}mapFrom +// HOST: call {{.*}} @.[[OFFLOAD_FROM]].region_id{{.*}} @[[MAPTYPES_FROM]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_FROM]](i{{[0-9]*}}* {{[^,]*}}%x) +// CHECK: call void ({{.*}}@[[OUTLINE_FROM:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_FROM]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) +void mapFrom() { + int x; + #pragma omp target teams firstprivate(x) map(from:x) + ; +} + +// HOST: define {{.*}}mapTo +// HOST: call {{.*}} @.[[OFFLOAD_TO]].region_id{{.*}} @[[MAPTYPES_TO]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_TO]](i{{[0-9]*}}* {{[^,]*}}%x) +// CHECK: call void ({{.*}}@[[OUTLINE_TO:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_TO]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) +void mapTo() { + int x; + #pragma omp target teams firstprivate(x) map(to:x) + ; +} + +// HOST: define {{.*}}mapAlloc +// HOST: call {{.*}} @.[[OFFLOAD_ALLOC]].region_id{{.*}} @[[MAPTYPES_ALLOC]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_ALLOC]](i{{[0-9]*}}* {{[^,]*}}%x) +// CHECK: call void ({{.*}}@[[OUTLINE_ALLOC:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_ALLOC]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) +void mapAlloc() { + int x; + #pragma omp target teams firstprivate(x) map(alloc:x) + ; +} + +// HOST: define {{.*}}mapArray +// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R0]].region_id{{.*}} @[[MAPTYPES_ARRAY_R0]] +// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R1]].region_id{{.*}} @[[MAPTYPES_ARRAY_R1]] +// +// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R0]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) +// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R0:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) +// CHECK: %.omp.reduction.red_list = alloca [1 x i8*] +// +// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R1]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) +// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R1:.omp_outlined.[.0-9]*]] +// +// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) +// CHECK: %.omp.reduction.red_list = alloca [1 x i8*] +void mapArray() { + int x[77], y[88], z[99]; + #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z) + ; + #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(to:x,y,z) + ; +} + +# if HAS_INT128 +// HOST-INT128: define {{.*}}mapInt128 +// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R0]].region_id{{.*}} @[[MAPTYPES_INT128_R0]] +// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R1]].region_id{{.*}} @[[MAPTYPES_INT128_R1]] +// +// INT128: define {{.*}} void @[[OFFLOAD_INT128_R0]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) +// INT128: call void ({{.*}}@[[OUTLINE_INT128_R0:.omp_outlined.[.0-9]*]] +// +// INT128: define {{.*}} void @[[OUTLINE_INT128_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) +// INT128: %.omp.reduction.red_list = alloca [1 x i8*] +// +// INT128: define {{.*}} void @[[OFFLOAD_INT128_R1]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) +// INT128: call void ({{.*}}@[[OUTLINE_INT128_R1:.omp_outlined.[.0-9]*]] +// +// INT128: define {{.*}} void @[[OUTLINE_INT128_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) +// INT128: %.omp.reduction.red_list = alloca [1 x i8*] +void mapInt128() { + __int128 x, y, z; + #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z) + ; + #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(from:x,y,z) + ; +} +# endif +#endif Index: clang/test/OpenMP/target_teams_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_map_messages.cpp +++ clang/test/OpenMP/target_teams_map_messages.cpp @@ -1,7 +1,10 @@ -// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized #ifdef CCODE void foo(int arg) { const int n = 0; @@ -536,10 +539,10 @@ #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}} foo(); -#pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as private}} +#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as private}} {} -#pragma omp target teams firstprivate(j) map(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}} +#pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}} {} #pragma omp target teams map(m)