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 @@ -9315,6 +9315,8 @@ "'ordered' clause with specified parameter">; def err_omp_expected_base_var_name : Error< "expected variable name as a base of the array %select{subscript|section}0">; +def err_omp_pointer_type_not_last : Error< + "only one level of indirection allowed in array section">; def err_omp_map_shared_storage : Error< "variable already marked as mapped in current construct">; def err_omp_invalid_map_type_for_directive : Error< 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 @@ -38,7 +38,7 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose); + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose); namespace { /// Default data sharing attributes, which can be applied to directive. @@ -2997,7 +2997,7 @@ if (isOpenMPTargetExecutionDirective(DKind)) { OMPClauseMappableExprCommon::MappableExprComponentList CurComponents; if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map, - /*NoDiagnose=*/true)) + Stack->getCurrentDirective(), /*NoDiagnose=*/true)) return; const auto *VD = cast( CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl()); @@ -14252,7 +14252,7 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose) { + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) { SourceLocation ELoc = E->getExprLoc(); SourceRange ERange = E->getSourceRange(); @@ -14277,6 +14277,10 @@ const Expr *RelevantExpr = nullptr; + // OpenMP 5.0 + // The target update construct was modified to allow array section that + // specify discontiguous storage. + // // 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. // @@ -14297,6 +14301,15 @@ bool AllowUnitySizeArraySection = true; bool AllowWholeSizeArraySection = true; + // OpenMP 5.0 [2.1.5, Array Sections] + // If a consecutive sequence of array subscript expressions appears in an + // array section, and the first subscript expression in the sequence uses + // the extended array section syntax defined in this section, then only the + // last subscript expression in the sequence may select array elements that + // have a pointer type. + bool IsLastArraySection = true; + bool HasMeetPointer = false; + while (!RelevantExpr) { E = E->IgnoreParenImpCasts(); @@ -14434,6 +14447,12 @@ return nullptr; } + if (HasMeetPointer && !IsLastArraySection) { + SemaRef.Diag(ELoc, diag::err_omp_pointer_type_not_last) + << CurE->getSourceRange(); + return nullptr; + } + bool NotWhole = checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, CurType); bool NotUnity = @@ -14448,7 +14467,8 @@ // pointer. Otherwise, only unitary sections are accepted. if (NotWhole || IsPointer) AllowWholeSizeArraySection = false; - } else if (AllowUnitySizeArraySection && NotUnity) { + } else if (DKind != OMPD_target_update && + (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( @@ -14481,6 +14501,9 @@ RelevantExpr = TE; } + IsLastArraySection = false; + if (IsPointer) HasMeetPointer = true; + // Record the component - we don't have any declaration associated. CurComponents.emplace_back(CurE, nullptr); } else { @@ -14952,7 +14975,8 @@ // Obtain the array or member expression bases if required. Also, fill the // components array with all the components identified in the process. const Expr *BE = checkMapClauseExpressionBase( - SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false); + SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(), + /*NoDiagnose=*/false); if (!BE) continue; diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -2695,7 +2695,7 @@ // CK19-DAG: [[SEC222222]] = load double***, double**** [[PTR]], // CK19: call void [[CALL42:@.+]](double*** {{[^,]+}}) - #pragma omp target map(mptras[:1][2][:13]) + #pragma omp target map(mptras[0][2][:13]) { mptras[1][2][3]++; } 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 @@ -258,19 +258,19 @@ #pragma omp target map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target map(mptr[:1][:2-1][2:4-3]) + #pragma omp target map(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target map(mptr[:1][:arg][2:4-3]) // correct if arg is 1. + #pragma omp target map(mptr[:1][:arg][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target map(mptr[:1][:2-1][0:2]) + #pragma omp target map(mptr[:1][:2-1][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} #pragma omp target map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} #pragma omp target map(r.ArrS[0].B) 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 @@ -205,19 +205,19 @@ #pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}} {} - #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target teams map(mptr[:1][:2-1][2:4-3]) + #pragma omp target teams map(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // correct if arg is 1. + #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target teams map(mptr[:1][:2-1][0:2]) + #pragma omp target teams map(mptr[:1][:2-1][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} - #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} #pragma omp target teams map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} {} - #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}} + #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{only one level of indirection allowed in array section}} {} #pragma omp target teams map(r.ArrS[0].B) diff --git a/clang/test/OpenMP/target_update_ast_print.cpp b/clang/test/OpenMP/target_update_ast_print.cpp --- a/clang/test/OpenMP/target_update_ast_print.cpp +++ b/clang/test/OpenMP/target_update_ast_print.cpp @@ -20,6 +20,40 @@ #pragma omp target update to(a) if(l>5) device(l) nowait depend(inout:l) #pragma omp target update from(b) if(l<5) device(l-1) nowait depend(inout:l) + + U marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) + +#pragma omp target update from(marr[2][0:2][0:2]) + +#pragma omp target update to(marr[:][0:2][0:2]) + +#pragma omp target update from(marr[:][0:2][0:2]) + +#pragma omp target update to(marr[:][:l][l:]) + +#pragma omp target update from(marr[:][:l][l:]) + +#pragma omp target update to(marr[:2][:1][:]) + +#pragma omp target update from(marr[:2][:1][:]) + +#pragma omp target update to(marr[:2][:][:1]) + +#pragma omp target update from(marr[:2][:][:1]) + +#pragma omp target update to(marr[:2][:][1:]) + +#pragma omp target update from(marr[:2][:][1:]) + +#pragma omp target update to(marr[:1][3:2][:2]) + +#pragma omp target update from(marr[:1][3:2][:2]) + +#pragma omp target update to(marr[:1][:2][0]) + +#pragma omp target update from(marr[:1][:2][0]) + return a + targ + (T)b; } // CHECK: static T a; @@ -37,6 +71,23 @@ // CHECK-NEXT: int l; // CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l) // CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l) +// CHECK-NEXT: marr[10][10][10]; +// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:][:l][l:]) +// CHECK-NEXT: #pragma omp target update from(marr[:][:l][l:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0]) int main(int argc, char **argv) { static int a; @@ -50,6 +101,37 @@ // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n) #pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n) // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n) + + float marr[10][10][10]; +// CHECK-NEXT: marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +#pragma omp target update from(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +#pragma omp target update to(marr[:][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:][0:2][0:2]) +#pragma omp target update from(marr[:][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:][0:2][0:2]) +#pragma omp target update to(marr[:][:n][n:]) +// CHECK-NEXT: #pragma omp target update to(marr[:][:n][n:]) +#pragma omp target update from(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:]) +#pragma omp target update to(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1]) +#pragma omp target update from(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1]) +#pragma omp target update to(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:]) +#pragma omp target update from(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:]) +#pragma omp target update to(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +#pragma omp target update from(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +#pragma omp target update to(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0]) +#pragma omp target update from(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0]) return foo(argc, f) + foo(argv[0][0], f) + a; } 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 @@ -80,6 +80,7 @@ const T (&l)[5] = da; T *m; S7 s7; + T ***mptr; #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'}} @@ -118,6 +119,9 @@ { #pragma omp target update from(s7.x) } +#pragma omp target update from(mptr[:2][2+2-4:1][0:5+5]) // expected-error 2 {{only one level of indirection allowed in array section}} +#pragma omp target update from(mptr[:1][:2-1][2:4-3]) // expected-error 2 {{only one level of indirection allowed in array section}} +#pragma omp target update from(mptr[:1][:2][0:2]) // expected-error 2 {{only one level of indirection allowed in array section}} return 0; } @@ -136,6 +140,7 @@ const int (&l)[5] = da; int *m; S7 s7; + int ***mptr; #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'}} @@ -172,6 +177,9 @@ { #pragma omp target update from(s7.x) } +#pragma omp target update from(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} return tmain(argc)+tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} expected-note {{in instantiation of function template specialization 'tmain' requested here}} } 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 @@ -80,6 +80,7 @@ T to; const T (&l)[5] = da; S7 s7; + T ***mptr; #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'}} @@ -118,6 +119,9 @@ { #pragma omp target update to(s7.x) } +#pragma omp target update to(mptr[:2][2+2-4:1][0:5+5]) // expected-error 2 {{only one level of indirection allowed in array section}} +#pragma omp target update to(mptr[:1][:2-1][2:4-3]) // expected-error 2 {{only one level of indirection allowed in array section}} +#pragma omp target update to(mptr[:1][:2][0:2]) // expected-error 2 {{only one level of indirection allowed in array section}} return 0; } @@ -135,6 +139,7 @@ const int (&l)[5] = da; S7 s7; int *m; + int ***mptr; #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'}} @@ -171,6 +176,9 @@ { #pragma omp target update to(s7.x) } +#pragma omp target update to(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} return tmain(argc)+tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} expected-note {{in instantiation of function template specialization 'tmain' requested here}} }