[OpenACC][CIR] Implement 'async'/'if' lowering for 'data' construct

These two are trivial, and work the same as the compute construct
versions of these, so this adds tests to do so, and adds them to the
implementation.

3 files changed, 119 insertions, 5 deletions

diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 47537d6..14c4532 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -204,7 +204,8 @@ public:
       if (!clause.getArchitectures().empty())
         operation.setDeviceType(
             decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
-    } else if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
+    } else if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp,
+                                      DataOp>) {
       // Nothing to do here, these constructs don't have any IR for these, as
       // they just modify the other clauses IR.  So setting of `lastDeviceType`
       // (done above) is all we need.
@@ -243,7 +244,7 @@ public:
   }
 
   void VisitAsyncClause(const OpenACCAsyncClause &clause) {
-    if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
+    if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>) {
       if (!clause.hasIntExpr()) {
         operation.setAsyncOnlyAttr(
             handleDeviceTypeAffectedClause(operation.getAsyncOnlyAttr()));
@@ -278,7 +279,7 @@ public:
 
   void VisitIfClause(const OpenACCIfClause &clause) {
     if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
-                               ShutdownOp, SetOp>) {
+                               ShutdownOp, SetOp, DataOp>) {
       operation.getIfCondMutable().append(
           createCondition(clause.getConditionExpr()));
     } else {
diff --git a/clang/test/CIR/CodeGenOpenACC/data.c b/clang/test/CIR/CodeGenOpenACC/data.c
index 9e636f6..29fd465 100644
--- a/clang/test/CIR/CodeGenOpenACC/data.c
+++ b/clang/test/CIR/CodeGenOpenACC/data.c
@@ -1,7 +1,9 @@
 // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
 
-void acc_data(void) {
-  // CHECK: cir.func @acc_data() {
+void acc_data(int cond) {
+  // CHECK: cir.func @acc_data(%[[ARG:.*]]: !s32i{{.*}}) {
+  // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+  // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
 
 #pragma acc data default(none)
   {
@@ -33,5 +35,79 @@ void acc_data(void) {
   // CHECK-NEXT: acc.terminator
   // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
 
+#pragma acc data default(none) async
+  {}
+  // CHECK-NEXT: acc.data {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>], defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) async(cond)
+  {}
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: acc.data async(%[[CONV_CAST]] : si32) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) async device_type(nvidia, radeon) async
+  {}
+  // CHECK-NEXT: acc.data {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>], defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) async(3) device_type(nvidia, radeon) async(cond)
+  {}
+  // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
+  // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: acc.data async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) async device_type(nvidia, radeon) async(cond)
+  {}
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: acc.data async(%[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<none>], defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) async(3) device_type(nvidia, radeon) async
+  {}
+  // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
+  // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
+  // CHECK-NEXT: acc.data async(%[[THREE_CAST]] : si32) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {asyncOnly = [#acc.device_type<nvidia>, #acc.device_type<radeon>], defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) if(cond)
+  {}
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+  // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) if(1)
+  {}
+  // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+  // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
+#pragma acc data default(none) if(cond == 1)
+  {}
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
+  // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1
+  // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {
+  // CHECK-NEXT: acc.terminator
+  // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+
   // CHECK-NEXT: cir.return
 }
diff --git a/clang/test/SemaOpenACC/data-construct-async-clause.c b/clang/test/SemaOpenACC/data-construct-async-clause.c
index 3c9fbae..6cb34b0 100644
--- a/clang/test/SemaOpenACC/data-construct-async-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-async-clause.c
@@ -31,4 +31,41 @@ void Test() {
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
 #pragma acc enter data copyin(I) async(I, I)
+  //
+  // expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async async
+  while(1);
+
+  // expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async(1) async(2)
+  while(1);
+
+  // expected-error@+2{{OpenACC 'async' clause cannot appear more than once on a 'data' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async(1) async(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async(1) device_type(*) async(1) async(2)
+  while(1);
+  // expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async device_type(*) async async
+  while(1);
+  // expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) async(1) device_type(*) async async(2)
+  while(1);
+
+  // expected-error@+3{{OpenACC 'async' clause cannot appear more than once in a 'device_type' region on a 'data' directive}}
+  // expected-note@+2{{previous clause is here}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc data default(none) device_type(*) async async
+  while(1);
 }