From 6cfad635d5aaa01abb82edc386329d8ed25078e1 Mon Sep 17 00:00:00 2001 From: erichkeane Date: Thu, 12 Dec 2024 09:41:05 -0800 Subject: [PATCH] [OpenACC] Implement 'if' clause sema for 'data' constructs This is another one that has no additional sema work other than enabling it, so this patch does just that. --- clang/include/clang/Basic/OpenACCKinds.h | 8 ++ clang/lib/Sema/SemaOpenACC.cpp | 14 +- .../AST/ast-print-openacc-data-construct.cpp | 16 +++ .../combined-construct-if-clause.c | 1 - .../SemaOpenACC/compute-construct-if-clause.c | 1 - .../SemaOpenACC/data-construct-if-ast.cpp | 128 ++++++++++++++++++ .../SemaOpenACC/data-construct-if-clause.c | 37 +++++ clang/test/SemaOpenACC/data-construct.cpp | 10 +- 8 files changed, 200 insertions(+), 15 deletions(-) create mode 100644 clang/test/SemaOpenACC/data-construct-if-ast.cpp create mode 100644 clang/test/SemaOpenACC/data-construct-if-clause.c diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h index ea0bf23468cb8b..7fb76271826a6b 100644 --- a/clang/include/clang/Basic/OpenACCKinds.h +++ b/clang/include/clang/Basic/OpenACCKinds.h @@ -158,6 +158,14 @@ inline bool isOpenACCCombinedDirectiveKind(OpenACCDirectiveKind K) { K == OpenACCDirectiveKind::KernelsLoop; } +// Tests 'K' to see if it is 'data', 'host_data', 'enter data', or 'exit data'. +inline bool isOpenACCDataDirectiveKind(OpenACCDirectiveKind K) { + return K == OpenACCDirectiveKind::Data || + K == OpenACCDirectiveKind::EnterData || + K == OpenACCDirectiveKind::ExitData || + K == OpenACCDirectiveKind::HostData; +} + enum class OpenACCAtomicKind : uint8_t { Read, Write, diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 981a98d024a281..ca073e93b4a869 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -641,16 +641,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on 'compute'/'combined' constructs, - // and 'compute'/'combined' constructs are the only construct that can do - // anything with this yet, so skip/treat as unimplemented in this case. + // Restrictions only properly implemented on 'compute'/'combined'/'data' + // constructs, and 'compute'/'combined'/'data' constructs are the only + // constructs that can do anything with this yet, so skip/treat as + // unimplemented in this case. if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) && - !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) + !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) && + !isOpenACCDataDirectiveKind(Clause.getDirectiveKind())) return isNotImplemented(); // There is no prose in the standard that says duplicates aren't allowed, // but this diagnostic is present in other compilers, as well as makes - // sense. + // sense. Prose DOES exist for 'data' and 'host_data', 'enter data' and 'exit + // data' both don't, but other implmementations do this. OpenACC issue 519 + // filed for the latter two. if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) return nullptr; diff --git a/clang/test/AST/ast-print-openacc-data-construct.cpp b/clang/test/AST/ast-print-openacc-data-construct.cpp index c65121317c099d..88540657dd52e3 100644 --- a/clang/test/AST/ast-print-openacc-data-construct.cpp +++ b/clang/test/AST/ast-print-openacc-data-construct.cpp @@ -27,4 +27,20 @@ void foo() { // CHECK-NOT: use_device(Var) #pragma acc host_data use_device(Var) ; + + int i; + int array[5]; + +// CHECK: #pragma acc data if(i == array[1]) +#pragma acc data default(none) if(i == array[1]) + ; +// CHECK: #pragma acc enter data if(i == array[1]) +#pragma acc enter data copyin(Var) if(i == array[1]) + ; +// CHECK: #pragma acc exit data if(i == array[1]) +#pragma acc exit data copyout(Var) if(i == array[1]) + ; +// CHECK: #pragma acc host_data if(i == array[1]) +#pragma acc host_data use_device(Var) if(i == array[1]) + ; } diff --git a/clang/test/SemaOpenACC/combined-construct-if-clause.c b/clang/test/SemaOpenACC/combined-construct-if-clause.c index c0069c4ee9ef44..09bd4dd190b6b2 100644 --- a/clang/test/SemaOpenACC/combined-construct-if-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-if-clause.c @@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) { #pragma acc kernels loop if (*I < *F) for (unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc data if (*I < *F) for (unsigned i = 0; i < 5; ++i); #pragma acc parallel loop if (*I < *F) diff --git a/clang/test/SemaOpenACC/compute-construct-if-clause.c b/clang/test/SemaOpenACC/compute-construct-if-clause.c index 1336bf2549f3ca..20d42a17cba141 100644 --- a/clang/test/SemaOpenACC/compute-construct-if-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-if-clause.c @@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) { #pragma acc kernels if (*I < *F) while(0); - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc data if (*I < *F) while(0); #pragma acc parallel loop if (*I < *F) diff --git a/clang/test/SemaOpenACC/data-construct-if-ast.cpp b/clang/test/SemaOpenACC/data-construct-if-ast.cpp new file mode 100644 index 00000000000000..5a7e5942211d38 --- /dev/null +++ b/clang/test/SemaOpenACC/data-construct-if-ast.cpp @@ -0,0 +1,128 @@ +// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s + +// Test this with PCH. +// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s +// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s + +#ifndef PCH_HELPER +#define PCH_HELPER +void NormalFunc(int j, float f) { + // CHECK: FunctionDecl{{.*}}NormalFunc + // CHECK-NEXT: ParmVarDecl + // CHECK-NEXT: ParmVarDecl + // CHECK-NEXT: CompoundStmt +#pragma acc data if( j < f) default(none) + ; + // CHECK-NEXT: OpenACCDataConstruct{{.*}}data + // CHECK-NEXT: if clause + // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' + // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'j' 'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' + // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float' + // CHECK-NEXT: NullStmt + +} + +int Global; + +template +void TemplFunc() { + // CHECK: FunctionTemplateDecl{{.*}}TemplFunc + // CHECK-NEXT: TemplateTypeParmDecl + + // Match the prototype: + // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc + // CHECK-NEXT: CompoundStmt + +#pragma acc data default(none) if(T::SomeFloat < typename T::IntTy{}) + ; + // CHECK-NEXT: OpenACCDataConstruct{{.*}}data + // CHECK-NEXT: if clause + // CHECK-NEXT: BinaryOperator{{.*}} '' '<' + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T' + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy' + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: NullStmt + +#pragma acc enter data copyin(Global) if(typename T::IntTy{}) + ; + // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}}enter data + // CHECK-NEXT: if clause + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename T::IntTy' 'typename T::IntTy' + // CHECK-NEXT: InitListExpr{{.*}} 'void' + // CHECK-NEXT: NullStmt + +#pragma acc exit data copyout(Global) if(T::SomeFloat) + ; + // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data + // CHECK-NEXT: if clause + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T' + // CHECK-NEXT: NullStmt + +#pragma acc host_data use_device(Global) if(T::BC) + ; + // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data + // CHECK-NEXT: if clause + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T' + // CHECK-NEXT: NullStmt + + // Match the instantiation: + // CHECK: FunctionDecl{{.*}}TemplFunc{{.*}}implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'InstTy' + // CHECK-NEXT: RecordType{{.*}} 'InstTy' + // CHECK-NEXT: CXXRecord{{.*}} 'InstTy' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCDataConstruct{{.*}}data + // CHECK-NEXT: if clause + // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' + // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float' + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy + // CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int' + // CHECK-NEXT: NullStmt + + // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}}enter data + // CHECK-NEXT: if clause + // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' + // CHECK-NEXT: CXXFunctionalCastExpr{{.*}}'typename InstTy::IntTy':'int' functional cast to typename struct InstTy::IntTy + // CHECK-NEXT: InitListExpr {{.*}}'typename InstTy::IntTy':'int' + // CHECK-NEXT: NullStmt + + // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data + // CHECK-NEXT: if clause + // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' + // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy' + // CHECK-NEXT: NullStmt + + // CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data + // CHECK-NEXT: if clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'bool' + // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'bool' + // CHECK-NEXT: MemberExpr{{.*}} .operator bool + // CHECK-NEXT: DeclRefExpr{{.*}} 'const BoolConversion' lvalue Var{{.*}} 'BC' 'const BoolConversion' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy' + // CHECK-NEXT: NullStmt + +} + +struct BoolConversion{ operator bool() const;}; +struct InstTy { + using IntTy = int; + static constexpr float SomeFloat = 5.0; + static constexpr BoolConversion BC; +}; + +void Instantiate() { + TemplFunc(); +} +#endif diff --git a/clang/test/SemaOpenACC/data-construct-if-clause.c b/clang/test/SemaOpenACC/data-construct-if-clause.c new file mode 100644 index 00000000000000..b57eb28d29dba2 --- /dev/null +++ b/clang/test/SemaOpenACC/data-construct-if-clause.c @@ -0,0 +1,37 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +void Foo() { + int Var; + // expected-warning@+1{{OpenACC clause 'default' not yet implemented}} +#pragma acc data default(present) if(1) + ; + // expected-warning@+3{{OpenACC clause 'default' not yet implemented}} + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'data' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc data default(present) if(1) if (2) + ; + + // expected-warning@+1{{OpenACC clause 'copyin' not yet implemented}} +#pragma acc enter data copyin(Var) if(1) + + // expected-warning@+3{{OpenACC clause 'copyin' not yet implemented}} + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'enter data' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc enter data copyin(Var) if(1) if (2) + + // expected-warning@+1{{OpenACC clause 'copyout' not yet implemented}} +#pragma acc exit data copyout(Var) if(1) + // expected-warning@+3{{OpenACC clause 'copyout' not yet implemented}} + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'exit data' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc exit data copyout(Var) if(1) if (2) + + // expected-warning@+1{{OpenACC clause 'use_device' not yet implemented}} +#pragma acc host_data use_device(Var) if(1) + ; + // expected-warning@+3{{OpenACC clause 'use_device' not yet implemented}} + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'host_data' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc host_data use_device(Var) if(1) if (2) + ; +} diff --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp index 2292065e781a2e..1b3c1985258e59 100644 --- a/clang/test/SemaOpenACC/data-construct.cpp +++ b/clang/test/SemaOpenACC/data-construct.cpp @@ -52,7 +52,6 @@ void AtLeastOneOf() { // OpenACC TODO: The following 'data' directives should diagnose, since they // don't have at least one of the above clauses. - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc data if(Var) ; @@ -80,7 +79,6 @@ void AtLeastOneOf() { // OpenACC TODO: The following 'enter data' directives should diagnose, since // they don't have at least one of the above clauses. - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc enter data if(Var) // expected-warning@+1{{OpenACC clause 'async' not yet implemented}} #pragma acc enter data async @@ -99,7 +97,6 @@ void AtLeastOneOf() { // OpenACC TODO: The following 'exit data' directives should diagnose, since // they don't have at least one of the above clauses. - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc exit data if(Var) // expected-warning@+1{{OpenACC clause 'async' not yet implemented}} #pragma acc exit data async @@ -116,7 +113,6 @@ void AtLeastOneOf() { // OpenACC TODO: The following 'host_data' directives should diagnose, since // they don't have at least one of the above clauses. - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} #pragma acc host_data if(Var) ; // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}} @@ -200,10 +196,8 @@ struct HasMembers { void HostDataRules() { int Var, Var2; - // TODO OpenACC: The following line should diagnose, since only 1 'if' is - // allowed per directive on host_data. - // expected-warning@+2{{OpenACC clause 'if' not yet implemented}} - // expected-warning@+1{{OpenACC clause 'if' not yet implemented}} + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'host_data' directive}} + // expected-note@+1{{previous clause is here}} #pragma acc host_data if(Var) if (Var2) ;