[PATCH] D69316: [OpenMP 5.0] target update list items need not be contiguous (Sema)

Chi Chun Chen via Phabricator via cfe-commits Wed, 30 Oct 2019 13:54:22 -0700

cchen updated this revision to Diff 227174.
cchen added a comment.

Rebase and fix add OpenMP 50 guard



Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D69316/new/

https://reviews.llvm.org/D69316

Files:
  clang/lib/Sema/SemaOpenMP.cpp
  clang/test/OpenMP/target_update_ast_print.cpp

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/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());
@@ -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;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

[PATCH] D69316: [OpenMP 5.0] target update list items need not be contiguous (Sema)

Reply via email to