[OpenMP] Extended parse for 'always' map modifier

[carlo.bertolli]

Summary

This patch allows the map modifier 'always' to be separated by the map type (to, from, tofrom) only by a whitespace, rather than strictly by a comma as in current trunk.

[ABataev]

LG

[carlo.bertolli]

Committed revision 302031.

lib/Parse/ParseOpenMP.cpp

if (IsMapClauseModifierToken(Tok)) {
           if (Data.MapType == OMPC_MAP_unknown ||
               Data.MapType == OMPC_MAP_always)
             Diag(Tok, diag::err_omp_unknown_map_type);
           ConsumeToken();
         } else {
           Data.MapType = OMPC_MAP_tofrom;
           Data.IsMapTypeImplicit = true;
         }
+      } else if (IsMapClauseModifierToken(PP.LookAhead(0))) {
+        if (PP.LookAhead(1).is(tok::colon)) {
+          Data.MapTypeModifier = Data.MapType;
+          if (Data.MapTypeModifier != OMPC_MAP_always) {
+            Diag(Tok, diag::err_omp_unknown_map_type_modifier);
+            Data.MapTypeModifier = OMPC_MAP_unknown;
+          } else
+            MapTypeModifierSpecified = true;
+
+          ConsumeToken();
+
+          Data.MapType =
+              IsMapClauseModifierToken(Tok)
+                  ? static_cast<OpenMPMapClauseKind>(
+                        getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)))
+                  : OMPC_MAP_unknown;
+          if (Data.MapType == OMPC_MAP_unknown ||
+              Data.MapType == OMPC_MAP_always)
+            Diag(Tok, diag::err_omp_unknown_map_type);
+          ConsumeToken();
+        } else {
+          Data.MapType = OMPC_MAP_tofrom;
+          Data.IsMapTypeImplicit = true;
+        }
       } else {
         Data.MapType = OMPC_MAP_tofrom;
         Data.IsMapTypeImplicit = true;
       }
     } else {
       Data.MapType = OMPC_MAP_tofrom;
       Data.IsMapTypeImplicit = true;
     }

test/OpenMP/target_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
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void foo() {}
 
 template <typename T, int C>
 T tmain(T argc, T *argv) {
-  T i, j, a[20];
+  T i, j, a[20], always;
 #pragma omp target
   foo();
 #pragma omp target if (target:argc > 0)
   foo();
 #pragma omp target if (C)
   foo();
 #pragma omp target map(i)
   foo();
 #pragma omp target map(a[0:10], i)
   foo();
 #pragma omp target map(to: i) map(from: j)
   foo();
 #pragma omp target map(always,alloc: i)
   foo();
+#pragma omp target map(always from: i)
+  foo();
+#pragma omp target map(always)
+  {always++;}
+#pragma omp target map(always,i)
+  {always++;i++;}
 #pragma omp target nowait
   foo();
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
   foo();
 #pragma omp target defaultmap(tofrom: scalar)
   foo();
   return 0;
 }

 // CHECK-NEXT: #pragma omp target map(tofrom: i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(always,alloc: i)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: always)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: always,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
 // CHECK-NEXT: int i, j, a[20]
 // CHECK-NEXT: #pragma omp target
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target if(target: argc > 0)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target if(5)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(tofrom: i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(always,alloc: i)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: always)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: always,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 // CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
 // CHECK-NEXT: char i, j, a[20]
 // CHECK-NEXT: #pragma omp target
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target if(target: argc > 0)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target if(1)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(tofrom: i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target map(always,alloc: i)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: always)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: always,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 
 // CHECK-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
-  int i, j, a[20];
+  int i, j, a[20], always;
 // CHECK-NEXT: int i, j, a[20]
 #pragma omp target
 // CHECK-NEXT: #pragma omp target
   foo();
 // CHECK-NEXT: foo();
 #pragma omp target if (argc > 0)
 // CHECK-NEXT: #pragma omp target if(argc > 0)
   foo();
// CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target map(always,alloc: i)
 // CHECK-NEXT: #pragma omp target map(always,alloc: i)
   foo();
 // CHECK-NEXT: foo();
 
+#pragma omp target map(always from: i)
+// CHECK-NEXT: #pragma omp target map(always,from: i)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(always)
+// CHECK-NEXT: #pragma omp target map(tofrom: always)
+  {always++;}
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: }
+
+#pragma omp target map(always,i)
+// CHECK-NEXT: #pragma omp target map(tofrom: always,i)
+  {always++;i++;}
+// CHECK-NEXT: {
+// CHECK-NEXT: always++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
+
 #pragma omp target nowait
 // CHECK-NEXT: #pragma omp target nowait
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
   foo();

test/OpenMP/target_map_messages.cpp

void func(int arg) {
     #pragma omp target map(always, tofrom: c[1:2],f)
     {}
     #pragma omp target map(always, tofrom: c,f[1:2])
     {}
     #pragma omp target map(always, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
     #pragma omp target map(always, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
+    #pragma omp target map(always)   // expected-error {{use of undeclared identifier 'always'}}
+    {}
     return;
   }
 };
 
 struct SB {
   unsigned A;
   unsigned B;
   float Arr[100];