[Clang][OpenMP] Add default map type for target enter/exit data

[doru1004]

Summary

Starting with OpenMP 5.2 the map type on target enter/exit data can be elided. When the map type is not specified it defaults to to for target enter data and to from for target exit data.

[ABataev]

clang/lib/Parse/ParseOpenMP.cpp

4367–4371

Add braces, otherwise the compiler generates a warning

[doru1004]

clang/lib/Parse/ParseOpenMP.cpp

4367–4371

Thanks for pointing it out!

[ABataev]

LG

[doru1004]

Commit: 9e595e911eb539caad99fd8642328007d47c6f4e

clang/lib/Parse/ParseOpenMP.cpp

if (ColonPresent) {
       IsInvalidMapperModifier = parseMapTypeModifiers(Data);
       if (!IsInvalidMapperModifier)
         parseMapType(*this, Data);
       else
         SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch);
     }
     if (Data.ExtraModifier == OMPC_MAP_unknown) {
       Data.ExtraModifier = OMPC_MAP_tofrom;
+      if (getLangOpts().OpenMP >= 52)
+        if (DKind == OMPD_target_enter_data)
+          Data.ExtraModifier = OMPC_MAP_to;
+        else if (DKind == OMPD_target_exit_data)
+          Data.ExtraModifier = OMPC_MAP_from;
       Data.IsMapTypeImplicit = true;
     }
 
     if (Tok.is(tok::colon))
       Data.ColonLoc = ConsumeToken();
   } else if (Kind == OMPC_to || Kind == OMPC_from) {
     while (Tok.is(tok::identifier)) {
       auto Modifier = static_cast<OpenMPMotionModifierKind>(

clang/lib/Sema/SemaOpenMP.cpp

for (Expr *RE : MVLI.VarList) {
     if (!checkTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
                            DSAS, Type, /*FullCheck=*/true))
       continue;
 
     if (CKind == OMPC_map) {
       // target enter data
       // OpenMP [2.10.2, Restrictions, p. 99]
       // A map-type must be specified in all map clauses and must be either
-      // to or alloc.
+      // to or alloc. Starting with OpenMP 5.2 the default map type is `to` if
+      // no map type is present.
       OpenMPDirectiveKind DKind = DSAS->getCurrentDirective();
       if (DKind == OMPD_target_enter_data &&
-          !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {
+          !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc ||
+            SemaRef.getLangOpts().OpenMP >= 52)) {
         SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
             << (IsMapTypeImplicit ? 1 : 0)
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // target exit_data
       // OpenMP [2.10.3, Restrictions, p. 102]
       // A map-type must be specified in all map clauses and must be either
-      // from, release, or delete.
+      // from, release, or delete. Starting with OpenMP 5.2 the default map
+      // type is `from` if no map type is present.
       if (DKind == OMPD_target_exit_data &&
           !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release ||
-            MapType == OMPC_MAP_delete)) {
+            MapType == OMPC_MAP_delete || SemaRef.getLangOpts().OpenMP >= 52)) {
         SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
             << (IsMapTypeImplicit ? 1 : 0)
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // The 'ompx_hold' modifier is specifically intended to be used on a

clang/test/OpenMP/target_enter_data_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
 // 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
 

clang/test/OpenMP/target_enter_data_ast_print_openmp52.cpp

+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck --check-prefix=CHECK --check-prefix=CHECK-52 %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+  T i_def, i;
+
+  i = argc;
+
+#pragma omp target enter data map(i_def)
+
+#pragma omp target enter data map(to: i)
+
+  return 0;
+}
+
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i_def){{$}}
+// CHECK-NEXT: #pragma omp target enter data map(to: i){{$}}
+
+// CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
+// CHECK-NEXT: int i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i_def){{$}}
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+
+// CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
+// CHECK-NEXT: char i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i_def){{$}}
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+
+int main (int argc, char **argv) {
+  int b_def, b;
+  static int a_def, a;
+// CHECK: static int a_def, a;
+
+#pragma omp target enter data map(a_def)
+// CHECK:      #pragma omp target enter data map(to: a_def)
+  a_def=2;
+// CHECK-NEXT: a_def = 2;
+
+#pragma omp target enter data map(to: a)
+// CHECK:      #pragma omp target enter data map(to: a)
+  a=2;
+// CHECK-NEXT: a = 2;
+
+#pragma omp target enter data map(b_def)
+// CHECK-NEXT:      #pragma omp target enter data map(to: b_def)
+
+#pragma omp target enter data map(to: b)
+// CHECK-NEXT:      #pragma omp target enter data map(to: b)
+
+  return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif

clang/test/OpenMP/target_exit_data_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=52 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
 // 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
 

clang/test/OpenMP/target_exit_data_ast_print_openmp52.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=52 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+  T i_def, i;
+
+  i = argc;
+#pragma omp target exit data map(i_def)
+
+#pragma omp target exit data map(from: i)
+
+  return 0;
+}
+
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i_def){{$}}
+// CHECK-NEXT: #pragma omp target exit data map(from: i){{$}}
+
+// CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
+// CHECK-NEXT: int i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i_def)
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+
+// CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
+// CHECK-NEXT: char i_def, i;
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i_def)
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+
+int main (int argc, char **argv) {
+  int b_def, b;
+  static int a_def, a;
+// CHECK: static int a_def, a;
+
+#pragma omp target exit data map(a_def)
+// CHECK:      #pragma omp target exit data map(from: a_def)
+  a_def=2;
+// CHECK-NEXT: a_def = 2;
+
+#pragma omp target exit data map(from: a)
+// CHECK:      #pragma omp target exit data map(from: a)
+  a=2;
+// CHECK-NEXT: a = 2;
+
+#pragma omp target exit data map(b_def)
+// CHECK-NEXT:      #pragma omp target exit data map(from: b_def)
+
+#pragma omp target exit data map(from: b)
+// CHECK-NEXT:      #pragma omp target exit data map(from: b)
+
+  return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif