[OpenACC] enable 'async' clause for combined constructs

No additional work required over what we did for other constructs, so
this is just adding the tests and enabling the clauses.

clang/lib/Sema/SemaOpenACC.cpp

@@ -798,10 +798,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitAsyncClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute' constructs, and
-  // 'compute' constructs are the only construct that can do anything with
-  // this yet, so skip/treat as unimplemented in this case.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+  // 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.
+  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
+      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
   // There is no prose in the standard that says duplicates aren't allowed,

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

@@ -1,6 +1,7 @@
 // RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
 
 void foo() {
+  int *iPtr;
 // CHECK: #pragma acc parallel loop
 // CHECK-NEXT: for (int i = 0; i < 5; ++i)
 // CHECK-NEXT: ;
@@ -98,11 +99,27 @@ void foo() {
   for(int i = 0;i<5;++i);
 
 // CHECK: #pragma acc parallel loop private(i, array[1], array, array[1:2])
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
 #pragma acc parallel loop private(i, array[1], array, array[1:2])
   for(int i = 0;i<5;++i);
 
 // CHECK: #pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
 #pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
   for(int i = 0;i<5;++i);
 
+  // CHECK: #pragma acc kernels loop async(*iPtr)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop async(*iPtr)
+  for(int i = 0;i<5;++i);
+
+  // CHECK: #pragma acc kernels loop async
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop async
+  for(int i = 0;i<5;++i);
+
 }

clang/test/SemaOpenACC/combined-construct-async-clause.c

@@ -0,0 +1,45 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+
+void Test() {
+#pragma acc parallel loop async
+  for (int i = 5; i < 10; ++i);
+#pragma acc parallel loop async(1)
+  for (int i = 5; i < 10; ++i);
+#pragma acc kernels loop async(1)
+  for (int i = 5; i < 10; ++i);
+#pragma acc kernels loop async(-51)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc serial loop async(1)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc serial loop async(1, 2)
+  for (int i = 5; i < 10; ++i);
+
+  struct NotConvertible{} NC;
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel loop async(NC)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc kernels loop async(getS())
+  for (int i = 5; i < 10; ++i);
+
+  struct Incomplete *SomeIncomplete;
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct Incomplete' invalid)}}
+#pragma acc kernels loop async(*SomeIncomplete)
+  for (int i = 5; i < 10; ++i);
+
+  enum E{A} SomeE;
+
+#pragma acc kernels loop async(SomeE)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}}
+#pragma acc loop async(1)
+  for(int i = 5; i < 10;++i);
+}

clang/test/SemaOpenACC/combined-construct-async-clause.cpp

@@ -0,0 +1,137 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NotConvertible{} NC;
+struct Incomplete *SomeIncomplete; // #INCOMPLETE
+enum E{} SomeE;
+enum class E2{} SomeE2;
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+struct ExplicitConvertOnly {
+  explicit operator int() const; // #EXPL_CONV
+} Explicit;
+
+struct AmbiguousConvert{
+  operator int(); // #AMBIG_INT
+  operator short(); // #AMBIG_SHORT
+  operator float();
+} Ambiguous;
+
+void Test() {
+#pragma acc parallel loop async
+  for (int i = 5; i < 10; ++i);
+#pragma acc parallel loop async(1)
+  for (int i = 5; i < 10; ++i);
+#pragma acc kernels loop async(-51)
+  for (int i = 5; i < 10; ++i);
+#pragma acc serial loop async(2)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel loop async(NC)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{OpenACC integer expression has incomplete class type 'struct Incomplete'}}
+  // expected-note@#INCOMPLETE{{forward declaration of 'Incomplete'}}
+#pragma acc kernels loop async(*SomeIncomplete)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop async(SomeE)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('enum E2' invalid}}
+#pragma acc kernels loop async(SomeE2)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop async(Convert)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels loop async(Explicit)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel loop async(Ambiguous)
+  for (int i = 5; i < 10; ++i);
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+  static constexpr AmbiguousConvert ACValue;
+  static constexpr ExplicitConvertOnly EXValue;
+
+  operator char();
+};
+
+template<typename T>
+void TestInst() {
+
+  // expected-error@+1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel loop async(HasInt::Invalid)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-note@#INST{{in instantiation of function template specialization 'TestInst<HasInt>' requested here}}
+#pragma acc kernels loop async(T::Invalid)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel loop async(HasInt::ACValue)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc kernels loop async(T::ACValue)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc parallel loop async(HasInt::EXValue)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels loop async(T::EXValue)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop async(HasInt::value)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc kernels loop async(T::value)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop async(HasInt::IntTy{})
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc kernels loop async(typename T::ShortTy{})
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop async(HasInt::IntTy{})
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc kernels loop async(typename T::ShortTy{})
+  for (int i = 5; i < 10; ++i);
+
+  HasInt HI{};
+  T MyT{};
+
+#pragma acc parallel loop async(HI)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc kernels loop async(MyT)
+  for (int i = 5; i < 10; ++i);
+}
+
+void Inst() {
+  TestInst<HasInt>(); // #INST
+}

clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c

@@ -190,8 +190,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto dtype(*)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop auto async
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
@@ -355,8 +353,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop dtype(*) auto
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop async auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
@@ -521,8 +517,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent dtype(*)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop independent async
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
@@ -686,8 +680,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop dtype(*) independent
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop async independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
@@ -858,8 +850,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq dtype(*)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop seq async
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}
@@ -1029,8 +1019,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop dtype(*) seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'async' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented}}
 #pragma acc parallel loop async seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}}

clang/test/SemaOpenACC/combined-construct-default-clause.cpp

@@ -7,28 +7,24 @@ void SingleOnly() {
 
   int i;
 
-  // expected-warning@+3{{OpenACC clause 'async' not yet implemented, clause ignored}}
   // expected-error@+2{{OpenACC 'default' clause cannot appear more than once on a 'parallel loop' directive}}
   // expected-note@+1{{previous clause is here}}
   #pragma acc parallel loop default(present) async default(none)
   for (unsigned I = 0; I < 5; ++I);
 
-  // expected-warning@+4{{OpenACC clause 'copy' not yet implemented, clause ignored}}
-  // expected-warning@+3{{OpenACC clause 'async' not yet implemented, clause ignored}}
+  // expected-warning@+3{{OpenACC clause 'copy' not yet implemented, clause ignored}}
   // expected-error@+2{{OpenACC 'default' clause cannot appear more than once on a 'serial loop' directive}}
   // expected-note@+1{{previous clause is here}}
   #pragma acc serial loop async default(present) copy(i) default(none) self
   for (unsigned I = 0; I < 5; ++I);
 
-  // expected-warning@+4{{OpenACC clause 'copy' not yet implemented, clause ignored}}
-  // expected-warning@+3{{OpenACC clause 'async' not yet implemented, clause ignored}}
+  // expected-warning@+3{{OpenACC clause 'copy' not yet implemented, clause ignored}}
   // expected-error@+2{{OpenACC 'default' clause cannot appear more than once on a 'kernels loop' directive}}
   // expected-note@+1{{previous clause is here}}
   #pragma acc kernels loop async default(present) copy(i) default(none) self
   for (unsigned I = 0; I < 5; ++I);
 
-  // expected-warning@+3{{OpenACC clause 'copy' not yet implemented, clause ignored}}
-  // expected-warning@+2{{OpenACC clause 'async' not yet implemented, clause ignored}}
+  // expected-warning@+2{{OpenACC clause 'copy' not yet implemented, clause ignored}}
   // expected-error@+1{{expected '('}}
   #pragma acc parallel loop async default(none) copy(i) default self
   for (unsigned I = 0; I < 5; ++I);

clang/test/SemaOpenACC/combined-construct-device_type-clause.c

@@ -207,7 +207,6 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc serial loop device_type(*) default_async(1)
   for(int i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel loop device_type(*) async
   for(int i = 0; i < 5; ++i);