Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -10888,10 +10888,11 @@ continue; } - // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] + // 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 - if (isOpenMPTargetExecutionDirective(CurrDir)) { + // attribute clause on the same construct unless the construct is a + // combined construct. + if (CurrDir == OMPD_target) { OpenMPClauseKind ConflictKind; if (DSAStack->checkMappableExprComponentListsForDecl( VD, /*CurrentRegionOnly=*/true, @@ -11123,10 +11124,11 @@ } } - // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] + // 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 - if (isOpenMPTargetExecutionDirective(CurrDir)) { + // attribute clause on the same construct unless the construct is a + // combined construct. + if (CurrDir == OMPD_target) { OpenMPClauseKind ConflictKind; if (DSAStack->checkMappableExprComponentListsForDecl( VD, /*CurrentRegionOnly=*/true, @@ -14225,10 +14227,11 @@ continue; } - // OpenMP 4.5 [2.15.5.1, Restrictions, p.3] + // 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 - if (VD && isOpenMPTargetExecutionDirective(DKind)) { + // attribute clause on the same construct unless the construct is a + // combined construct. + if (VD && 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) @@ -15193,8 +15196,9 @@ // Check if the declaration in the clause does not show up in any data // sharing attribute. + OpenMPDirectiveKind CurrDir = DSAStack->getCurrentDirective(); DSAStackTy::DSAVarData DVar = DSAStack->getTopDSA(D, /*FromParent=*/false); - if (isOpenMPPrivate(DVar.CKind)) { + if (CurrDir == OMPD_target && isOpenMPPrivate(DVar.CKind)) { Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa) << getOpenMPClauseName(DVar.CKind) << getOpenMPClauseName(OMPC_is_device_ptr) Index: clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp @@ -87,10 +87,10 @@ #pragma omp target parallel for is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int ii=0; ii<10; ii++) ; -#pragma omp target parallel for firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for' directive}} expected-note{{defined as firstprivate}} +#pragma omp target parallel for firstprivate(ps) is_device_ptr(ps) for (int ii=0; ii<10; ii++) ; -#pragma omp target parallel for private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for' directive}} expected-note{{defined as private}} +#pragma omp target parallel for private(ps) is_device_ptr(ps) for (int ii=0; ii<10; ii++) ; Index: clang/test/OpenMP/target_parallel_for_simd_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_simd_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_parallel_for_simd_is_device_ptr_messages.cpp @@ -228,16 +228,16 @@ #pragma omp target parallel for simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} +#pragma omp target parallel for simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target parallel for simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} +#pragma omp target parallel for simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} expected-note{{defined as private}} +#pragma omp target parallel for simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return 0; @@ -323,16 +323,16 @@ #pragma omp target parallel for simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} +#pragma omp target parallel for simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target parallel for simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} +#pragma omp target parallel for simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target parallel for simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel for simd' directive}} expected-note{{defined as private}} +#pragma omp target parallel for simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} Index: clang/test/OpenMP/target_parallel_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_parallel_is_device_ptr_messages.cpp @@ -189,13 +189,13 @@ {} #pragma omp target parallel is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} {} -#pragma omp target parallel is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} +#pragma omp target parallel is_device_ptr(ps) firstprivate(ps) {} -#pragma omp target parallel firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} expected-note{{defined as firstprivate}} +#pragma omp target parallel firstprivate(ps) is_device_ptr(ps) {} -#pragma omp target parallel is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} +#pragma omp target parallel is_device_ptr(ps) private(ps) {} -#pragma omp target parallel private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} expected-note{{defined as private}} +#pragma omp target parallel private(ps) is_device_ptr(ps) {} return 0; } @@ -258,13 +258,13 @@ {} #pragma omp target parallel is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} {} -#pragma omp target parallel is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} +#pragma omp target parallel is_device_ptr(ps) firstprivate(ps) {} -#pragma omp target parallel firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} expected-note{{defined as firstprivate}} +#pragma omp target parallel firstprivate(ps) is_device_ptr(ps) {} -#pragma omp target parallel is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} +#pragma omp target parallel is_device_ptr(ps) private(ps) {} -#pragma omp target parallel private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target parallel' directive}} expected-note{{defined as private}} +#pragma omp target parallel private(ps) is_device_ptr(ps) {} return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} } Index: clang/test/OpenMP/target_teams_distribute_parallel_for_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_is_device_ptr_messages.cpp @@ -228,16 +228,16 @@ #pragma omp target teams distribute parallel for is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} +#pragma omp target teams distribute parallel for is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute parallel for firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} +#pragma omp target teams distribute parallel for is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute parallel for private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return 0; @@ -323,16 +323,16 @@ #pragma omp target teams distribute parallel for is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} +#pragma omp target teams distribute parallel for is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute parallel for firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} +#pragma omp target teams distribute parallel for is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute parallel for private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} 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 @@ -240,7 +240,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) 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_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_is_device_ptr_messages.cpp @@ -228,16 +228,16 @@ #pragma omp target teams distribute parallel for simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} +#pragma omp target teams distribute parallel for simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute parallel for simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} +#pragma omp target teams distribute parallel for simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute parallel for simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return 0; @@ -323,16 +323,16 @@ #pragma omp target teams distribute parallel for simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} +#pragma omp target teams distribute parallel for simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute parallel for simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} +#pragma omp target teams distribute parallel for simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute parallel for simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute parallel for simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' 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 @@ -240,7 +240,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) 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 @@ -147,7 +147,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) for (i = 0; i < argc; ++i) foo(); return 0; Index: clang/test/OpenMP/target_teams_distribute_simd_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_is_device_ptr_messages.cpp @@ -228,16 +228,16 @@ #pragma omp target teams distribute simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} +#pragma omp target teams distribute simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} +#pragma omp target teams distribute simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return 0; @@ -323,16 +323,16 @@ #pragma omp target teams distribute simd is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} +#pragma omp target teams distribute simd is_device_ptr(ps) firstprivate(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams distribute simd firstprivate(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} +#pragma omp target teams distribute simd is_device_ptr(ps) private(ps) for (int i=0; i<100; i++) ; -#pragma omp target teams distribute simd private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams distribute simd' directive}} expected-note{{defined as private}} +#pragma omp target teams distribute simd private(ps) is_device_ptr(ps) for (int i=0; i<100; i++) ; return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} 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 @@ -240,7 +240,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) 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 @@ -132,7 +132,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) for (int k = 0; k < argc; ++k) ++k; return 0; Index: clang/test/OpenMP/target_teams_is_device_ptr_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_is_device_ptr_messages.cpp +++ clang/test/OpenMP/target_teams_is_device_ptr_messages.cpp @@ -189,13 +189,13 @@ {} #pragma omp target teams is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} {} -#pragma omp target teams is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} +#pragma omp target teams is_device_ptr(ps) firstprivate(ps) {} -#pragma omp target teams firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams firstprivate(ps) is_device_ptr(ps) {} -#pragma omp target teams is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} +#pragma omp target teams is_device_ptr(ps) private(ps) {} -#pragma omp target teams private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as private}} +#pragma omp target teams private(ps) is_device_ptr(ps) {} return 0; } @@ -258,13 +258,13 @@ {} #pragma omp target teams is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}} {} -#pragma omp target teams is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} +#pragma omp target teams is_device_ptr(ps) firstprivate(ps) {} -#pragma omp target teams firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as firstprivate}} +#pragma omp target teams firstprivate(ps) is_device_ptr(ps) {} -#pragma omp target teams is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} +#pragma omp target teams is_device_ptr(ps) private(ps) {} -#pragma omp target teams private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as private}} +#pragma omp target teams private(ps) is_device_ptr(ps) {} return tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} } Index: clang/test/OpenMP/target_teams_map_codegen.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_teams_map_codegen.cpp @@ -0,0 +1,64 @@ +// Test host codegen. +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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=45 -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 + +// Test target codegen - host bc file has to be created first. +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp -fopenmp-version=45 -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 +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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=45 -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=45 -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=45 -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 + +// 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: define {{.*}} void @[[OFFLOAD_PRIVATE]]() +// CHECK-NOT: {{^}$}} +// 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) + ; +} + +// CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}} %x, i{{[0-9]*}} %y) +// CHECK-NOT: {{^}$}} +// 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) + ; +} + +// CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}} %x, i{{[0-9]*}}* {{[^,]*}} %y) +// CHECK-NOT: {{^}$}} +// CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]] + +// CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}} %x, i{{[0-9]*}}* {{[^,]*}} %y) +// CHECK-NOT: {{^}$}} +// 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) + ; +} +#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 @@ -536,10 +536,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) {} -#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) {} #pragma omp target teams map(m)