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()); @@ -14723,7 +14723,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(); @@ -14748,6 +14748,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. // @@ -14919,7 +14923,9 @@ // pointer. Otherwise, only unitary sections are accepted. if (NotWhole || IsPointer) AllowWholeSizeArraySection = false; - } else if (AllowUnitySizeArraySection && NotUnity) { + } else if (SemaRef.getLangOpts().OpenMP >= 50 && + 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( @@ -15431,7 +15437,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_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; }