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@+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@+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@+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@+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@+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@+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@+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@+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@+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@+2{{OpenACC 'alwaysout' modifier not valid on 'copyin' clause}} // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copyin' clause}} @@ -90,5 +91,6 @@ void ModList() { #pragma acc enter data copyin(alwaysout: V1) // expected-error@+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@+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@+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@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}} // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}} @@ -90,6 +91,7 @@ void ModList() { #pragma acc exit data copyout(alwaysout: V1) // expected-error@+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@+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@+3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} // expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}} // expected-error@+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@+1{{OpenACC 'always' modifier not valid on 'create' clause}} #pragma acc data create(always: V1) // expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} @@ -88,12 +88,14 @@ void ModList() { // expected-error@+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@+4{{OpenACC 'always' modifier not valid on 'create' clause}} - // expected-error@+3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} - // expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}} - // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'create' clause}} -#pragma acc enter data create(always, alwaysin, alwaysout, zero, readonly: V1) + // expected-error@+5{{OpenACC 'always' modifier not valid on 'create' clause}} + // expected-error@+4{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} + // expected-error@+3{{OpenACC 'alwaysout' modifier not valid on 'create' clause}} + // expected-error@+2{{OpenACC 'readonly' modifier not valid on 'create' clause}} + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'create' clause}} +#pragma acc enter data create(always, alwaysin, alwaysout, zero, readonly, capture: V1) // expected-error@+1{{OpenACC 'always' modifier not valid on 'create' clause}} #pragma acc enter data create(always: V1) // expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} @@ -102,5 +104,8 @@ void ModList() { #pragma acc enter data create(alwaysout: V1) // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'create' clause}} #pragma acc enter data create(readonly: V1) + // expected-error@+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@+2{{OpenACC 'readonly' modifier not valid on 'copy' clause}} // expected-error@+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@+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@+2{{OpenACC 'alwaysout' modifier not valid on 'copyin' clause}} // expected-error@+1{{OpenACC 'zero' modifier not valid on 'copyin' clause}} @@ -325,7 +326,10 @@ void ModList() { #pragma acc declare copyin(alwaysout: V6) // expected-error@+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@+1{{OpenACC 'capture' modifier not valid on 'copyin' clause}} +#pragma acc declare copyin(capture: V7B) + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'copyin' clause}} +#pragma acc declare copyin(always, alwaysin, readonly, capture: V8) // expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}} // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}} @@ -334,13 +338,17 @@ void ModList() { #pragma acc declare copyout(alwaysout: V10) // expected-error@+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@+4{{OpenACC 'always' modifier not valid on 'create' clause}} - // expected-error@+3{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} - // expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'create' clause}} - // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'create' clause}} -#pragma acc declare create(always, alwaysin, alwaysout, zero, readonly: V13) + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'copyout' clause}} +#pragma acc declare copyout(capture: V11B) + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'copyout' clause}} +#pragma acc declare copyout(always, alwaysin, zero, capture: V12) + + // expected-error@+5{{OpenACC 'always' modifier not valid on 'create' clause}} + // expected-error@+4{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} + // expected-error@+3{{OpenACC 'alwaysout' modifier not valid on 'create' clause}} + // expected-error@+2{{OpenACC 'readonly' modifier not valid on 'create' clause}} + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'create' clause}} +#pragma acc declare create(always, alwaysin, alwaysout, zero, readonly, capture: V13) // expected-error@+1{{OpenACC 'always' modifier not valid on 'create' clause}} #pragma acc declare create(always: V14) // expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'create' clause}} @@ -349,5 +357,8 @@ void ModList() { #pragma acc declare create(alwaysout: V16) // expected-error@+1{{OpenACC 'readonly' modifier not valid on 'create' clause}} #pragma acc declare create(readonly: V17) + #pragma acc declare create(zero: V18) + // expected-error@+1{{OpenACC 'capture' modifier not valid on 'create' clause}} +#pragma acc declare create(capture: V19) } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
