[OpenACC] Implement 'capture' modifier Sema/AST

The 'capture' modifier is an OpenACC 3.3NEXT (AKA 3.4) feature, which
permits a new kind of identifying the memory location of variables in a
data clause. However, it is only valid on data, combined, or compute
constructs.

This patch implements all of the proper restrictions.

Author: erichkeane

clang/include/clang/Basic/OpenACCKinds.h

@@ -642,7 +642,8 @@ enum class OpenACCModifierKind : uint8_t {
   AlwaysOut = 1 << 2,
   Readonly = 1 << 3,
   Zero = 1 << 4,
-  LLVM_MARK_AS_BITMASK_ENUM(Zero)
+  Capture = 1 << 5,
+  LLVM_MARK_AS_BITMASK_ENUM(Capture)
 };
 
 inline bool isOpenACCModifierBitSet(OpenACCModifierKind List,
@@ -690,6 +691,13 @@ inline StreamTy &printOpenACCModifierKind(StreamTy &Out,
     Out << "zero";
     First = false;
   }
+
+  if (isOpenACCModifierBitSet(Mods, OpenACCModifierKind::Capture)) {
+    if (!First)
+      Out << ", ";
+    Out << "capture";
+    First = false;
+  }
   return Out;
 }
 inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,

clang/lib/Parse/ParseOpenACC.cpp

@@ -691,6 +691,7 @@ OpenACCModifierKind Parser::tryParseModifierList(OpenACCClauseKind CK) {
         .Case("alwaysout", OpenACCModifierKind::AlwaysOut)
         .Case("readonly", OpenACCModifierKind::Readonly)
         .Case("zero", OpenACCModifierKind::Zero)
+        .Case("capture", OpenACCModifierKind::Capture)
         .Default(OpenACCModifierKind::Invalid);
   };
 

clang/lib/Sema/SemaOpenACCClause.cpp

@@ -198,31 +198,50 @@ class SemaOpenACCClauseVisitor {
       Mods = CheckSingle(Mods, ValidKinds, OpenACCModifierKind::AlwaysOut);
       Mods = CheckSingle(Mods, ValidKinds, OpenACCModifierKind::Readonly);
       Mods = CheckSingle(Mods, ValidKinds, OpenACCModifierKind::Zero);
+      Mods = CheckSingle(Mods, ValidKinds, OpenACCModifierKind::Capture);
       return Mods;
     };
 
+    // The 'capture' modifier is only valid on copyin, copyout, and create on
+    // structured data or compute constructs (which also includes combined).
+    bool IsStructuredDataOrCompute =
+        Clause.getDirectiveKind() == OpenACCDirectiveKind::Data ||
+        isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) ||
+        isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind());
+
     switch (Clause.getClauseKind()) {
     default:
       llvm_unreachable("Only for copy, copyin, copyout, create");
     case OpenACCClauseKind::Copy:
     case OpenACCClauseKind::PCopy:
     case OpenACCClauseKind::PresentOrCopy:
+      // COPY: Capture always
       return Check(OpenACCModifierKind::Always | OpenACCModifierKind::AlwaysIn |
-                   OpenACCModifierKind::AlwaysOut);
+                   OpenACCModifierKind::AlwaysOut |
+                   OpenACCModifierKind::Capture);
     case OpenACCClauseKind::CopyIn:
     case OpenACCClauseKind::PCopyIn:
     case OpenACCClauseKind::PresentOrCopyIn:
+      // COPYIN: Capture only struct.data & compute
       return Check(OpenACCModifierKind::Always | OpenACCModifierKind::AlwaysIn |
-                   OpenACCModifierKind::Readonly);
+                   OpenACCModifierKind::Readonly |
+                   (IsStructuredDataOrCompute ? OpenACCModifierKind::Capture
+                                              : OpenACCModifierKind::Invalid));
     case OpenACCClauseKind::CopyOut:
     case OpenACCClauseKind::PCopyOut:
     case OpenACCClauseKind::PresentOrCopyOut:
+      // COPYOUT: Capture only struct.data & compute
       return Check(OpenACCModifierKind::Always | OpenACCModifierKind::AlwaysIn |
-                   OpenACCModifierKind::Zero);
+                   OpenACCModifierKind::Zero |
+                   (IsStructuredDataOrCompute ? OpenACCModifierKind::Capture
+                                              : OpenACCModifierKind::Invalid));
     case OpenACCClauseKind::Create:
     case OpenACCClauseKind::PCreate:
     case OpenACCClauseKind::PresentOrCreate:
-      return Check(OpenACCModifierKind::Zero);
+      // CREATE: Capture only struct.data & compute
+      return Check(OpenACCModifierKind::Zero |
+                   (IsStructuredDataOrCompute ? OpenACCModifierKind::Capture
+                                              : OpenACCModifierKind::Invalid));
     }
     llvm_unreachable("didn't return from switch above?");
   }

clang/test/SemaOpenACC/combined-construct-copy-clause.c

@@ -82,6 +82,8 @@ void ModList() {
   // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copy' clause}}
 #pragma acc kernels loop copy(zero: V1)
   for(int i = 5; i < 10;++i);
-#pragma acc parallel loop copy(always, alwaysin, alwaysout: V1)
+#pragma acc parallel loop copy(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel loop copy(always, alwaysin, alwaysout, capture: V1)
   for(int i = 5; i < 10;++i);
 }

clang/test/SemaOpenACC/combined-construct-copyin-clause.c

@@ -89,6 +89,8 @@ void ModList() {
   // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
 #pragma acc kernels loop copyin(zero: V1)
   for(int i = 5; i < 10;++i);
-#pragma acc parallel loop copyin(always, alwaysin, readonly: V1)
+#pragma acc parallel loop copyin(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel loop copyin(always, alwaysin, readonly, capture: V1)
   for(int i = 5; i < 10;++i);
 }

clang/test/SemaOpenACC/combined-construct-copyout-clause.c

@@ -88,7 +88,9 @@ void ModList() {
   // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
 #pragma acc kernels loop copyout(readonly: V1)
   for(int i = 0; i < 6;++i);
-#pragma acc parallel loop copyout(always, alwaysin, zero: V1)
-  for(int i = 0; i < 6;++i);
+#pragma acc parallel loop copyout(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel loop copyout(always, alwaysin, zero, capture: V1)
+  for(int i = 5; i < 10;++i);
 }
 

clang/test/SemaOpenACC/combined-construct-create-clause.c

@@ -99,4 +99,6 @@ void ModList() {
   for(int i = 5; i < 10;++i);
 #pragma acc kernels loop create(zero: V1)
   for(int i = 5; i < 10;++i);
+#pragma acc parallel loop create(capture:V1)
+  for(int i = 5; i < 10;++i);
 }

clang/test/SemaOpenACC/compute-construct-copy-clause.c

@@ -95,6 +95,8 @@ void ModList() {
   // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copy' clause}}
 #pragma acc kernels copy(zero: V1)
   for(int i = 5; i < 10;++i);
-#pragma acc parallel copy(always, alwaysin, alwaysout: V1)
+#pragma acc parallel copy(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel copy(always, alwaysin, alwaysout, capture: V1)
   for(int i = 5; i < 10;++i);
 }

clang/test/SemaOpenACC/compute-construct-copyin-clause.c

@@ -88,6 +88,8 @@ void ModList() {
   // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
 #pragma acc kernels copyin(zero: V1)
   for(int i = 5; i < 10;++i);
-#pragma acc parallel copyin(always, alwaysin, readonly: V1)
+#pragma acc parallel copyin(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel copyin(always, alwaysin, readonly, capture: V1)
   for(int i = 5; i < 10;++i);
 }

clang/test/SemaOpenACC/compute-construct-copyout-clause.c

@@ -88,6 +88,8 @@ void ModList() {
   // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
 #pragma acc kernels copyout(readonly: V1)
   for(int i = 0; i < 6;++i);
-#pragma acc parallel copyout(always, alwaysin, zero: V1)
-  for(int i = 0; i < 6;++i);
+#pragma acc parallel copyout(capture:V1)
+  for(int i = 5; i < 10;++i);
+#pragma acc parallel copyout(always, alwaysin, zero, capture: V1)
+  for(int i = 5; i < 10;++i);
 }