caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[OpenACC] Implement 'capture' modifier Sema/AST

Browse Source 
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.
This commit is contained in:
erichkeane
2025-05-29 13:08:18 -07:00
parent 1a1927abd3
commit c66dbbe385
16 changed files with 99 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
clang/include/clang/Basic/OpenACCKinds.h
View File

@@ -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,

1
clang/lib/Parse/ParseOpenACC.cpp
View File

@@ -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);
};

27
clang/lib/Sema/SemaOpenACCClause.cpp
View File

@@ -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?");
}

4
clang/test/SemaOpenACC/combined-construct-copy-clause.c
View File

@@ -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);
}

4
clang/test/SemaOpenACC/combined-construct-copyin-clause.c
View File

@@ -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);
}

6
clang/test/SemaOpenACC/combined-construct-copyout-clause.c
View File

@@ -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);
}

2
clang/test/SemaOpenACC/combined-construct-create-clause.c
View File

@@ -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);
}

4
clang/test/SemaOpenACC/compute-construct-copy-clause.c
View File

@@ -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);
}

4
clang/test/SemaOpenACC/compute-construct-copyin-clause.c
View File

@@ -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);
}

6
clang/test/SemaOpenACC/compute-construct-copyout-clause.c
View File

@@ -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);
}

2
clang/test/SemaOpenACC/compute-construct-create-clause.c
View File

@@ -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);
}

4
clang/test/SemaOpenACC/data-construct-copy-clause.c
View File

@@ -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)
;
}

6
clang/test/SemaOpenACC/data-construct-copyin-clause.c
View File

@@ -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)
}

6
clang/test/SemaOpenACC/data-construct-copyout-clause.c
View File

@@ -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)
}

17
clang/test/SemaOpenACC/data-construct-create-clause.c
View File

@@ -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)
}

31
clang/test/SemaOpenACC/declare-construct.cpp
View File

@@ -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@+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@+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@+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)
}