cchen created this revision.
cchen added reviewers: ABataev, sfantao.
Herald added subscribers: cfe-commits, guansong.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.
- Removed a restriction that list items for target update
directive must have contiguous storage.
- A restriction was added for array sections that array of pointers
is still disallowed. (also modify a test case in target_map_codegen
that is not valid anymore in OpenMP 5.0)
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D69316
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_map_codegen.cpp
clang/test/OpenMP/target_map_messages.cpp
clang/test/OpenMP/target_teams_map_messages.cpp
clang/test/OpenMP/target_update_ast_print.cpp
clang/test/OpenMP/target_update_from_messages.cpp
clang/test/OpenMP/target_update_to_messages.cpp
Index: clang/test/OpenMP/target_update_to_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_to_messages.cpp
+++ 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<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
}
Index: clang/test/OpenMP/target_update_from_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_from_messages.cpp
+++ 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<int, 3>(argc)+tmain<to, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
}
Index: clang/test/OpenMP/target_update_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_update_ast_print.cpp
+++ 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;
}
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
@@ -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)
Index: clang/test/OpenMP/target_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_map_messages.cpp
+++ 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)
Index: clang/test/OpenMP/target_map_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen.cpp
+++ 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]++;
}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ 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<ValueDecl>(
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;
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ 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<
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
