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

[OpenACC] Fix 'copyout' allowed modifiers alwaysin vs alwaysout

Browse Source 
While doing lowering, I discovered that the restriction onthe allowed
modifiers for 'copyout' didn't make sense! After discussion on the
OpenACC standards mailing list I discovered that this was a copy/paste
error during standardization that they intend to fix, and really meant
for copyout to allow alwaysout instead of alwaysin.

When implementing, I blindly followed the standard :)

This patch corrects the implementation to do what was meant.
This commit is contained in:
erichkeane
2025-06-26 10:24:06 -07:00
parent 0b8a656ba1
commit 0a2b6f6c1c
9 changed files with 38 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -231,8 +231,8 @@ class SemaOpenACCClauseVisitor {
case OpenACCClauseKind::PCopyOut:
case OpenACCClauseKind::PresentOrCopyOut:
// COPYOUT: Capture only struct.data & compute
return Check(OpenACCModifierKind::Always | OpenACCModifierKind::AlwaysIn |
OpenACCModifierKind::Zero |
return Check(OpenACCModifierKind::Always |
OpenACCModifierKind::AlwaysOut | OpenACCModifierKind::Zero |
(IsStructuredDataOrCompute ? OpenACCModifierKind::Capture
: OpenACCModifierKind::Invalid));
case OpenACCClauseKind::Create:

4
clang/test/AST/ast-print-openacc-combined-construct.cpp
View File

@@ -179,8 +179,8 @@ void foo() {
#pragma acc parallel loop copyin(i, array[1], array, array[1:2]) pcopyin(readonly:i, array[1], array, array[1:2]) present_or_copyin(always, alwaysin: i, array[1], array, array[1:2])
for(int i = 0;i<5;++i);
// CHECK: #pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysin: i, array[1], array, array[1:2])
#pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysin: i, array[1], array, array[1:2])
// CHECK: #pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysout: i, array[1], array, array[1:2])
#pragma acc parallel loop copyout(i, array[1], array, array[1:2]) pcopyout(zero: i, array[1], array, array[1:2]) present_or_copyout(always, alwaysout: i, array[1], array, array[1:2])
for(int i = 0;i<5;++i);
// CHECK: #pragma acc parallel loop create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])

10
clang/test/SemaOpenACC/combined-construct-copyout-ast.cpp
View File

@@ -14,7 +14,7 @@ void NormalUses(float *PointerParam) {
// CHECK: ParmVarDecl
// CHECK-NEXT: CompoundStmt
#pragma acc parallel loop copyout(GlobalArray) pcopyout(zero:PointerParam[Global]) present_or_copyout(always, alwaysin: Global)
#pragma acc parallel loop copyout(GlobalArray) pcopyout(zero:PointerParam[Global]) present_or_copyout(always, alwaysout: Global)
for (unsigned i = 0; i < 5; ++i);
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: copyout clause
@@ -25,7 +25,7 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}}'PointerParam' 'float *'
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysout
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
// CHECK-NEXT: For
// CHECK: NullStmt
@@ -42,10 +42,10 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: ParmVarDecl{{.*}} referenced u 'U'
// CHECK-NEXT: CompoundStmt
#pragma acc parallel loop copyout(always, alwaysin: t) pcopyout(zero: NTTP, u) present_or_copyout(u[0:t])
#pragma acc parallel loop copyout(always, alwaysout: t) pcopyout(zero: NTTP, u) present_or_copyout(u[0:t])
for (unsigned i = 0; i < 5; ++i);
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: copyout clause modifiers: always, alwaysin
// CHECK-NEXT: copyout clause modifiers: always, alwaysout
// CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 't' 'T'
// CHECK-NEXT: pcopyout clause modifiers: zero
// CHECK-NEXT: DeclRefExpr{{.*}}'auto' lvalue NonTypeTemplateParm{{.*}} 'NTTP' 'auto &'
@@ -73,7 +73,7 @@ void TemplUses(T t, U u) {
// #pragma acc parallel copyout(t) pcopyout(zero: NTTP, u) present_or_copyout(u[0:t])
// CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
// CHECK-NEXT: copyout clause modifiers: always, alwaysin
// CHECK-NEXT: copyout clause modifiers: always, alwaysout
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
// CHECK-NEXT: pcopyout clause modifiers: zero
// CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'const unsigned int' lvalue

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

@@ -78,19 +78,19 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
}
void ModList() {
int V1;
// expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc parallel loop copyout(always, alwaysin, alwaysout, zero, readonly: V1)
for(int i = 0; i < 6;++i);
// expected-error@+1{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
#pragma acc serial loop copyout(alwaysout: V1)
// expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
#pragma acc serial loop copyout(alwaysin: V1)
for(int i = 0; i < 6;++i);
// 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(capture:V1)
for(int i = 5; i < 10;++i);
#pragma acc parallel loop copyout(always, alwaysin, zero, capture: V1)
#pragma acc parallel loop copyout(always, alwaysout, zero, capture: V1)
for(int i = 5; i < 10;++i);
}

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

@@ -78,18 +78,18 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
}
void ModList() {
int V1;
// expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc parallel copyout(always, alwaysin, alwaysout, zero, readonly: V1)
for(int i = 0; i < 6;++i);
// expected-error@+1{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
#pragma acc serial copyout(alwaysout: V1)
// expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
#pragma acc serial copyout(alwaysin: V1)
for(int i = 0; i < 6;++i);
// 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(capture:V1)
for(int i = 5; i < 10;++i);
#pragma acc parallel copyout(always, alwaysin, zero, capture: V1)
#pragma acc parallel copyout(always, alwaysout, zero, capture: V1)
for(int i = 5; i < 10;++i);
}

12
clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
View File

@@ -133,7 +133,7 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
#pragma acc parallel copyout(GlobalArray) pcopyout(zero:PointerParam[Global]) present_or_copyout(always, alwaysin: Global)
#pragma acc parallel copyout(GlobalArray) pcopyout(zero:PointerParam[Global]) present_or_copyout(always, alwaysout: Global)
while(true);
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
// CHECK-NEXT: copyout clause
@@ -144,7 +144,7 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}}'PointerParam' 'float *'
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysout
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
@@ -363,7 +363,7 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
#pragma acc parallel copyout(t) pcopyout(zero:NTTP, u) present_or_copyout(always, alwaysin: u[0:t])
#pragma acc parallel copyout(t) pcopyout(zero:NTTP, u) present_or_copyout(always, alwaysout: u[0:t])
while(true);
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
// CHECK-NEXT: copyout clause
@@ -371,7 +371,7 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: pcopyout clause modifiers: zero
// CHECK-NEXT: DeclRefExpr{{.*}}'auto' lvalue NonTypeTemplateParm{{.*}} 'NTTP' 'auto &'
// CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysout
// CHECK-NEXT: ArraySectionExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
@@ -569,7 +569,7 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
// #pragma acc parallel copyout(t) pcopyout(zero:NTTP, u) present_or_copyout(always, alwaysin: u[0:t])
// #pragma acc parallel copyout(t) pcopyout(zero:NTTP, u) present_or_copyout(always, alwaysout: u[0:t])
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
// CHECK-NEXT: copyout clause
// CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
@@ -578,7 +578,7 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'auto &' depth 0 index 0 NTTP
// CHECK-NEXT: DeclRefExpr{{.*}}'const unsigned int' lvalue Var{{.*}} 'CEVar' 'const unsigned int'
// CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: always, alwaysout
// CHECK-NEXT: ArraySectionExpr
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
// CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'

6
clang/test/SemaOpenACC/data-construct-copyout-ast.cpp
View File

@@ -54,7 +54,7 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: ParmVarDecl{{.*}} referenced u 'U'
// CHECK-NEXT: CompoundStmt
#pragma acc data copyout(t) pcopyout(zero: NTTP, u) present_or_copyout(alwaysin: u[0:t])
#pragma acc data copyout(t) pcopyout(zero: NTTP, u) present_or_copyout(alwaysout: u[0:t])
;
// CHECK-NEXT: OpenACCDataConstruct{{.*}} data
// CHECK-NEXT: copyout clause
@@ -62,7 +62,7 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: pcopyout clause modifiers: zero
// CHECK-NEXT: DeclRefExpr{{.*}}'auto' lvalue NonTypeTemplateParm{{.*}} 'NTTP' 'auto &'
// CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
// CHECK-NEXT: present_or_copyout clause modifiers: alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: alwaysout
// CHECK-NEXT: ArraySectionExpr
// CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
@@ -103,7 +103,7 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'auto &' depth 0 index 0 NTTP
// CHECK-NEXT: DeclRefExpr{{.*}}'const unsigned int' lvalue Var{{.*}} 'CEVar' 'const unsigned int'
// CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'
// CHECK-NEXT: present_or_copyout clause modifiers: alwaysin
// CHECK-NEXT: present_or_copyout clause modifiers: alwaysout
// CHECK-NEXT: ArraySectionExpr
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'int *' <LValueToRValue>
// CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'

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

@@ -74,20 +74,22 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
void ModList() {
int V1;
// expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc data copyout(always, alwaysin, alwaysout, zero, readonly: V1)
// expected-error@+1{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
#pragma acc data copyout(alwaysin: V1)
#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(capture: V1)
#pragma acc data copyout(always, alwaysin, zero, capture: V1)
#pragma acc data copyout(always, alwaysout, zero, capture: V1)
// expected-error@+2{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc exit data copyout(always, alwaysin, alwaysout, zero, readonly: V1)
// expected-error@+1{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
#pragma acc exit data copyout(alwaysin: V1)
#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)

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

@@ -331,17 +331,18 @@ void ModList() {
// 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@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(always, alwaysin, alwaysout, zero, readonly: V9)
// expected-error@+1{{OpenACC 'alwaysout' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(alwaysout: V10)
// expected-error@+1{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(alwaysin: V10)
// expected-error@+1{{OpenACC 'readonly' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(readonly: V11)
// expected-error@+1{{OpenACC 'capture' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(capture: V11B)
// expected-error@+2{{OpenACC 'alwaysin' modifier not valid on 'copyout' clause}}
// expected-error@+1{{OpenACC 'capture' modifier not valid on 'copyout' clause}}
#pragma acc declare copyout(always, alwaysin, zero, capture: V12)
#pragma acc declare copyout(always, alwaysin, alwaysout, 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}}