[clang] c66dbbe - [OpenACC] Implement 'capture' modifier Sema/AST

via cfe-commits cfe-commits at lists.llvm.org
Fri May 30 09:31:30 PDT 2025


Author: erichkeane
Date: 2025-05-30T09:31:25-07:00
New Revision: c66dbbe385561a349ec854e545bbab395ec6dcf8

URL: https://github.com/llvm/llvm-project/commit/c66dbbe385561a349ec854e545bbab395ec6dcf8
DIFF: https://github.com/llvm/llvm-project/commit/c66dbbe385561a349ec854e545bbab395ec6dcf8.diff

LOG: [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.

Added: 
    

Modified: 
    clang/include/clang/Basic/OpenACCKinds.h
    clang/lib/Parse/ParseOpenACC.cpp
    clang/lib/Sema/SemaOpenACCClause.cpp
    clang/test/SemaOpenACC/combined-construct-copy-clause.c
    clang/test/SemaOpenACC/combined-construct-copyin-clause.c
    clang/test/SemaOpenACC/combined-construct-copyout-clause.c
    clang/test/SemaOpenACC/combined-construct-create-clause.c
    clang/test/SemaOpenACC/compute-construct-copy-clause.c
    clang/test/SemaOpenACC/compute-construct-copyin-clause.c
    clang/test/SemaOpenACC/compute-construct-copyout-clause.c
    clang/test/SemaOpenACC/compute-construct-create-clause.c
    clang/test/SemaOpenACC/data-construct-copy-clause.c
    clang/test/SemaOpenACC/data-construct-copyin-clause.c
    clang/test/SemaOpenACC/data-construct-copyout-clause.c
    clang/test/SemaOpenACC/data-construct-create-clause.c
    clang/test/SemaOpenACC/declare-construct.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index a798401dfa9f5..a7a29e2add20a 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/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,

diff  --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 3539278c7ff65..ca4f878464c4f 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/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);
   };
 

diff  --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 7249602f9cb04..8ee27694b7376 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/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?");
   }

diff  --git a/clang/test/SemaOpenACC/combined-construct-copy-clause.c b/clang/test/SemaOpenACC/combined-construct-copy-clause.c
index 07c412c621ff1..42aa8a62f47cc 100644
--- a/clang/test/SemaOpenACC/combined-construct-copy-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-copy-clause.c
@@ -82,6 +82,8 @@ void ModList() {
   // expected-error at +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);
 }

diff  --git a/clang/test/SemaOpenACC/combined-construct-copyin-clause.c b/clang/test/SemaOpenACC/combined-construct-copyin-clause.c
index b4a6eafdb9ebd..d8778e53c784d 100644
--- a/clang/test/SemaOpenACC/combined-construct-copyin-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-copyin-clause.c
@@ -89,6 +89,8 @@ void ModList() {
   // expected-error at +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);
 }

diff  --git a/clang/test/SemaOpenACC/combined-construct-copyout-clause.c b/clang/test/SemaOpenACC/combined-construct-copyout-clause.c
index 6621adb5c6124..2ca70736d30e3 100644
--- a/clang/test/SemaOpenACC/combined-construct-copyout-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-copyout-clause.c
@@ -88,7 +88,9 @@ void ModList() {
   // expected-error at +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);
 }
 

diff  --git a/clang/test/SemaOpenACC/combined-construct-create-clause.c b/clang/test/SemaOpenACC/combined-construct-create-clause.c
index bf7dfe83a0511..82d00bb110f71 100644
--- a/clang/test/SemaOpenACC/combined-construct-create-clause.c
+++ b/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);
 }

diff  --git a/clang/test/SemaOpenACC/compute-construct-copy-clause.c b/clang/test/SemaOpenACC/compute-construct-copy-clause.c
index 67682488fbe89..5e50eb6140231 100644
--- a/clang/test/SemaOpenACC/compute-construct-copy-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copy-clause.c
@@ -95,6 +95,8 @@ void ModList() {
   // expected-error at +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);
 }

diff  --git a/clang/test/SemaOpenACC/compute-construct-copyin-clause.c b/clang/test/SemaOpenACC/compute-construct-copyin-clause.c
index eaa8a604df32a..81a80fd68ab54 100644
--- a/clang/test/SemaOpenACC/compute-construct-copyin-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copyin-clause.c
@@ -88,6 +88,8 @@ void ModList() {
   // expected-error at +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);
 }

diff  --git a/clang/test/SemaOpenACC/compute-construct-copyout-clause.c b/clang/test/SemaOpenACC/compute-construct-copyout-clause.c
index f1ea21d0824cc..34f8d5aa60e09 100644
--- a/clang/test/SemaOpenACC/compute-construct-copyout-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copyout-clause.c
@@ -88,6 +88,8 @@ void ModList() {
   // expected-error at +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);
 }

diff  --git a/clang/test/SemaOpenACC/compute-construct-create-clause.c b/clang/test/SemaOpenACC/compute-construct-create-clause.c
index 926c5b88a5115..236d504f7c845 100644
--- a/clang/test/SemaOpenACC/compute-construct-create-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-create-clause.c
@@ -97,6 +97,6 @@ void ModList() {
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
 #pragma acc serial create(readonly: V1)
   for(int i = 5; i < 10;++i);
-#pragma acc kernels loop create(zero: V1)
+#pragma acc parallel create(capture:V1)
   for(int i = 5; i < 10;++i);
 }

diff  --git a/clang/test/SemaOpenACC/data-construct-copy-clause.c b/clang/test/SemaOpenACC/data-construct-copy-clause.c
index b19f3f465caaf..724d73f9d14de 100644
--- a/clang/test/SemaOpenACC/data-construct-copy-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-copy-clause.c
@@ -79,7 +79,9 @@ void ModList() {
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copy' clause}}
 #pragma acc data copy(zero: V1)
   ;
-#pragma acc data copy(always, alwaysin, alwaysout: V1)
+#pragma acc data copy(capture:V1)
+  ;
+#pragma acc data copy(always, alwaysin, alwaysout, capture: V1)
 ;
 }
 

diff  --git a/clang/test/SemaOpenACC/data-construct-copyin-clause.c b/clang/test/SemaOpenACC/data-construct-copyin-clause.c
index 9b51a28d51a3d..0b1bd3abed53d 100644
--- a/clang/test/SemaOpenACC/data-construct-copyin-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-copyin-clause.c
@@ -81,7 +81,8 @@ void ModList() {
 #pragma acc data copyin(alwaysout: V1)
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
 #pragma acc data copyin(zero: V1)
-#pragma acc data copyin(always, alwaysin, readonly: V1)
+#pragma acc data copyin(capture: V1)
+#pragma acc data copyin(always, alwaysin, readonly, capture: V1)
 
   // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'copyin' clause}}
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
@@ -90,5 +91,6 @@ void ModList() {
 #pragma acc enter data copyin(alwaysout: V1)
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
 #pragma acc enter data copyin(zero: V1)
-#pragma acc enter data copyin(always, alwaysin, readonly: V1)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyin' clause}}
+#pragma acc enter data copyin(capture: V1)
 }

diff  --git a/clang/test/SemaOpenACC/data-construct-copyout-clause.c b/clang/test/SemaOpenACC/data-construct-copyout-clause.c
index 0c2264c56fcb8..deaf48fc77597 100644
--- a/clang/test/SemaOpenACC/data-construct-copyout-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-copyout-clause.c
@@ -81,7 +81,8 @@ void ModList() {
 #pragma acc data copyout(alwaysout: V1)
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
 #pragma acc data copyout(readonly: V1)
-#pragma acc data copyout(always, alwaysin, zero: V1)
+#pragma acc data copyout(capture: V1)
+#pragma acc data copyout(always, alwaysin, zero, capture: V1)
 
   // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
@@ -90,6 +91,7 @@ void ModList() {
 #pragma acc exit data copyout(alwaysout: V1)
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
 #pragma acc exit data copyout(readonly: V1)
-#pragma acc exit data copyout(always, alwaysin, zero: V1)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyout' clause}}
+#pragma acc exit data copyout(capture: V1)
 }
 

diff  --git a/clang/test/SemaOpenACC/data-construct-create-clause.c b/clang/test/SemaOpenACC/data-construct-create-clause.c
index 560c6b65cc502..b81125d2a7dc2 100644
--- a/clang/test/SemaOpenACC/data-construct-create-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-create-clause.c
@@ -78,7 +78,7 @@ void ModList() {
   // expected-error at +3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
   // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}}
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
-#pragma acc data create(always, alwaysin, alwaysout, zero, readonly: V1)
+#pragma acc data create(always, alwaysin, alwaysout, zero, readonly, capture: V1)
   // expected-error at +1{{OpenACC 'always' modifier not valid on 'create' clause}}
 #pragma acc data create(always: V1)
   // expected-error at +1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
@@ -88,12 +88,14 @@ void ModList() {
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
 #pragma acc data create(readonly: V1)
 #pragma acc data create(zero: V1)
+#pragma acc data create(zero, capture: V1)
 
-  // expected-error at +4{{OpenACC 'always' modifier not valid on 'create' clause}}
-  // expected-error at +3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
-  // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}}
-  // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
-#pragma acc enter data create(always, alwaysin, alwaysout, zero, readonly: V1)
+  // expected-error at +5{{OpenACC 'always' modifier not valid on 'create' clause}}
+  // expected-error at +4{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
+  // expected-error at +3{{OpenACC 'alwaysout' modifier not valid on 'create' clause}}
+  // expected-error at +2{{OpenACC 'readonly' modifier not valid on 'create' clause}}
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'create' clause}}
+#pragma acc enter data create(always, alwaysin, alwaysout, zero, readonly, capture: V1)
   // expected-error at +1{{OpenACC 'always' modifier not valid on 'create' clause}}
 #pragma acc enter data create(always: V1)
   // expected-error at +1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
@@ -102,5 +104,8 @@ void ModList() {
 #pragma acc enter data create(alwaysout: V1)
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
 #pragma acc enter data create(readonly: V1)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'create' clause}}
+#pragma acc enter data create(capture: V1)
+
 #pragma acc enter data create(zero: V1)
 }

diff  --git a/clang/test/SemaOpenACC/declare-construct.cpp b/clang/test/SemaOpenACC/declare-construct.cpp
index 25038b5bf242c..ff8a2ec64c31d 100644
--- a/clang/test/SemaOpenACC/declare-construct.cpp
+++ b/clang/test/SemaOpenACC/declare-construct.cpp
@@ -307,8 +307,8 @@ struct Struct2 {
 };
 
 void ModList() {
-  int V1, V2, V3, V4, V5, V6, V7, V8, V9, V10,
-      V11, V12, V13, V14, V15, V16, V17, V18;
+  int V1, V2, V3, V4, V4B, V5, V6, V7, V7B, V8, V9, V10,
+      V11, V11B, V12, V13, V14, V15, V16, V17, V18, V19;
   // expected-error at +2{{OpenACC 'readonly' modifier not valid on 'copy' clause}}
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copy' clause}}
 #pragma acc declare copy(always, alwaysin, alwaysout, zero, readonly: V1)
@@ -316,7 +316,8 @@ void ModList() {
 #pragma acc declare copy(readonly: V2)
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copy' clause}}
 #pragma acc declare copy(zero: V3)
-#pragma acc declare copy(always, alwaysin, alwaysout: V4)
+#pragma acc declare copy(capture: V4)
+#pragma acc declare copy(always, alwaysin, alwaysout, capture: V4B)
 
   // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'copyin' clause}}
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
@@ -325,7 +326,10 @@ void ModList() {
 #pragma acc declare copyin(alwaysout: V6)
   // expected-error at +1{{OpenACC 'zero' modifier not valid on 'copyin' clause}}
 #pragma acc declare copyin(zero: V7)
-#pragma acc declare copyin(always, alwaysin, readonly: V8)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyin' clause}}
+#pragma acc declare copyin(capture: V7B)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyin' clause}}
+#pragma acc declare copyin(always, alwaysin, readonly, capture: V8)
 
   // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
@@ -334,13 +338,17 @@ void ModList() {
 #pragma acc declare copyout(alwaysout: V10)
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
 #pragma acc declare copyout(readonly: V11)
-#pragma acc declare copyout(always, alwaysin, zero: V12)
-
-  // expected-error at +4{{OpenACC 'always' modifier not valid on 'create' clause}}
-  // expected-error at +3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
-  // expected-error at +2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}}
-  // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
-#pragma acc declare create(always, alwaysin, alwaysout, zero, readonly: V13)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyout' clause}}
+#pragma acc declare copyout(capture: V11B)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'copyout' clause}}
+#pragma acc declare copyout(always, alwaysin, zero, capture: V12)
+
+  // expected-error at +5{{OpenACC 'always' modifier not valid on 'create' clause}}
+  // expected-error at +4{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
+  // expected-error at +3{{OpenACC 'alwaysout' modifier not valid on 'create' clause}}
+  // expected-error at +2{{OpenACC 'readonly' modifier not valid on 'create' clause}}
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'create' clause}}
+#pragma acc declare create(always, alwaysin, alwaysout, zero, readonly, capture: V13)
   // expected-error at +1{{OpenACC 'always' modifier not valid on 'create' clause}}
 #pragma acc declare create(always: V14)
   // expected-error at +1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}}
@@ -349,5 +357,8 @@ void ModList() {
 #pragma acc declare create(alwaysout: V16)
   // expected-error at +1{{OpenACC 'readonly' modifier not valid on 'create' clause}}
 #pragma acc declare create(readonly: V17)
+
 #pragma acc declare create(zero: V18)
+  // expected-error at +1{{OpenACC 'capture' modifier not valid on 'create' clause}}
+#pragma acc declare create(capture: V19)
 }


        


More information about the cfe-commits mailing list