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

[cchen]

Summary

Removed a restriction that list items for target update directive must have contiguous storage.

[cchen]

Fix summary message

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

14912

Ad the check for the version of OpenMP standard, it must be allowed only for OpenMP >= 50.

14926–14929

Same here, allow only for OpenMP >= 50. Also, split the patch into 2 parts, one for pointers stuff and another one for target update directive.

[cchen]

Remove the restriction rule for array section to another patch

[cchen]

Rebase and fix add OpenMP 50 guard

[ABataev]

What about codegen tests?

clang/test/OpenMP/target_update_ast_print.cpp

33

It seems to me, the tests are not updated. All non-contiguous tests must pass only if OpenMP version is 50

[cchen]

Add and fix test.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

14926–14929

SemaRef.getLangOpts().OpenMP < 50 || DKind != OMPD_target_update

clang/test/OpenMP/target_update_codegen.cpp

325–326

We still map the whole array here, right?

clang/lib/Sema/SemaOpenMP.cpp

 
 //===----------------------------------------------------------------------===//
 // Stack of data-sharing attributes for variables
 //===----------------------------------------------------------------------===//
 
 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.
 enum DefaultDataSharingAttributes {
   DSA_unspecified = 0, /// Data sharing attribute not specified.
   DSA_none = 1 << 0,   /// Default data sharing attribute 'none'.
   DSA_shared = 1 << 1, /// Default data sharing attribute 'shared'.
 };
if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParens())) {
         if (DVar.CKind != OMPC_unknown)
           ImplicitFirstprivate.push_back(E);
       }
       return;
     }
     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());
       if (!Stack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&CurComponents](
                   OMPClauseMappableExprCommon::MappableExprComponentListRef
                       StackComponents,

 
 // Return the expression of the base of the mappable expression or null if it
 // cannot be determined and do all the necessary checks to see if the expression
 // is valid as a standalone mappable expression. In the process, record all the
 // components of the expression.
 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();
 
   // The base of elements of list in a map clause have to be either:
   //  - a reference to variable or field.
   //  - a member expression.
   //  - an array expression.
   //
   // E.g. if we have the expression 'r.S.Arr[:12]', we want to retrieve the
   // reference to 'r'.
   //
   // If we have:
   //
   // struct SS {
   //   Bla S;
   //   foo() {
   //     #pragma omp target map (S.Arr[:12]);
   //   }
   // }
   //
   // We want to retrieve the member expression 'this->S';
 
   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.
   //
   // For this restriction it is sufficient that we make sure only references
   // to variables or fields and array expressions, and that no array sections
   // exist except in the rightmost expression (unless they cover the whole
   // dimension of the array). E.g. these would be invalid:
   //
if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
       bool IsPointer = CurType->isAnyPointerType();
 
       if (!IsPointer && !CurType->isArrayType()) {
         SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
             << 0 << CurE->getSourceRange();
         return nullptr;
       }
 
       bool NotWhole =
           checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, CurType);
       bool NotUnity =
           checkArrayExpressionDoesNotReferToUnitySize(SemaRef, CurE, CurType);
 
       if (AllowWholeSizeArraySection) {
         // Any array section is currently allowed. Allowing a whole size array
         // section implies allowing a unity array section as well.
         //
         // If this array section refers to the whole dimension we can still
         // accept other array sections before this one, except if the base is a
         // 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) ||
+                  SemaRef.getLangOpts().OpenMP < 50) &&
+                 (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(
             ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
             << CurE->getSourceRange();
         return nullptr;
       }
 
for (Expr *RE : MVLI.VarList) {
     }
 
     OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
     ValueDecl *CurDeclaration = nullptr;
 
     // 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;
 
     assert(!CurComponents.empty() &&
            "Invalid mappable expression information.");
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) {
       // Add store "this" pointer to class in DSAStackTy for future checking

clang/test/OpenMP/target_update_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
+
+// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void foo() {}
 
 template <class T, class U>
 T foo(T targ, U uarg) {
   static T a;
   U b;
   int l;
 #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)
-  return a + targ + (T)b;
-}
+
 // CHECK:      static T a;
 // CHECK-NEXT: U b;
 // 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:      static int a;
 // CHECK-NEXT: float b;
 // 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:      static char a;
 // CHECK-NEXT: float b;
 // 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)
 
+#ifdef OMP5
+  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])
+
+// OMP5: marr[10][10][10];
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][:l][l:])
+// OMP5-NEXT: #pragma omp target update from(marr[:][:l][l:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+#endif
+
+  return a + targ + (T)b;
+}
+
 int main(int argc, char **argv) {
   static int a;
   int n;
   float f;
 
 // CHECK:      static int a;
 // CHECK-NEXT: int n;
 // CHECK-NEXT: float f;
 #pragma omp target update to(a) if(f>0.0) device(n) nowait depend(in:n)
 // 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)
+
+#ifdef OMP5
+  float marr[10][10][10];
+// OMP5: marr[10][10][10];
+#pragma omp target update to(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+#pragma omp target update from(marr[2][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+#pragma omp target update to(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+#pragma omp target update from(marr[:][0:2][0:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+#pragma omp target update to(marr[:][:n][n:])
+// OMP5: #pragma omp target update to(marr[:][:n][n:])
+#pragma omp target update from(marr[:2][:1][:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:])
+#pragma omp target update to(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1])
+#pragma omp target update from(marr[:2][:][:1])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1])
+#pragma omp target update to(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:])
+#pragma omp target update from(marr[:2][:][1:])
+// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:])
+#pragma omp target update to(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+#pragma omp target update from(marr[:1][3:2][:2])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+#pragma omp target update to(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0])
+#pragma omp target update from(marr[:1][:2][0])
+// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0])
+#endif
+
   return foo(argc, f) + foo(argv[0][0], f) + a;
 }
 
 #endif

clang/test/OpenMP/target_update_codegen.cpp

void device_side_scan(int arg) {
   // CK4: tgt_target_data_update
   // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK4: ret
   // TCK4-NOT: tgt_target_data_update
   #pragma omp target update from(arg) if(arg) device(4)
   {++arg;}
 }
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
+// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-64
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
+// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
+
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK5
+
+// CK5: [[SIZE:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 80]
+// CK5: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+void target_update_contiguous() {
+  int marr[10][10][10];
+
+  // CK5-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [10 x [10 x [10 x i32]]]**
+  // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [10 x i32]**
+  // CK5-DAG: store [10 x [10 x [10 x i32]]]* %marr, [10 x [10 x [10 x i32]]]** [[BPC0]]
+  // CK5-DAG: store [10 x i32]* {{.+}}, [10 x i32]** [[PC0]]
+  #pragma omp target update to(marr[2][0:2][0:2])
+}
+#endif
 #endif