[OpenACC][CIR] Implement lowering for 'set' clauses

The 'set' clauses are default_async, device_num, and if.  The latter two
are implemented identically to other constructs by that name.

default_async works exactly like device_num, (and others) that take an
int-expression.

Author: erichkeane

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

@@ -170,6 +170,7 @@ class OpenACCClauseCIREmitter final
         break;
       }
     } else {
+      // Combined Constructs left.
       return clauseNotImplemented(clause);
     }
   }
@@ -208,6 +209,7 @@ class OpenACCClauseCIREmitter final
       // they just modify the other clauses IR.  So setting of `lastDeviceType`
       // (done above) is all we need.
     } else {
+      // update, data, loop, routine, combined remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -221,6 +223,7 @@ class OpenACCClauseCIREmitter final
     } else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
       llvm_unreachable("num_workers not valid on serial");
     } else {
+      // Combined Remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -234,6 +237,7 @@ class OpenACCClauseCIREmitter final
     } else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
       llvm_unreachable("vector_length not valid on serial");
     } else {
+      // Combined remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -250,6 +254,7 @@ class OpenACCClauseCIREmitter final
             createIntExpr(clause.getIntExpr()), &range));
       }
     } else {
+      // Data, enter data, exit data, update, wait, combined remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -266,19 +271,21 @@ class OpenACCClauseCIREmitter final
         llvm_unreachable("var-list version of self shouldn't get here");
       }
     } else {
+      // update and combined remain.
       return clauseNotImplemented(clause);
     }
   }
 
   void VisitIfClause(const OpenACCIfClause &clause) {
     if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
-                               ShutdownOp>) {
+                               ShutdownOp, SetOp>) {
       operation.getIfCondMutable().append(
           createCondition(clause.getConditionExpr()));
     } else {
       // 'if' applies to most of the constructs, but hold off on lowering them
       // until we can write tests/know what we're doing with codegen to make
       // sure we get it right.
+      // Enter data, exit data, host_data, update, wait, combined remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -287,8 +294,23 @@ class OpenACCClauseCIREmitter final
     if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
       operation.getDeviceNumOperandMutable().append(
           createIntExpr(clause.getIntExpr()));
+    } else if constexpr (isOneOfTypes<OpTy, SetOp>) {
+      // This is only a separate case because the getter name is different in
+      // 'set' for some reason.
+      operation.getDeviceNumMutable().append(
+          createIntExpr(clause.getIntExpr()));
     } else {
-      return clauseNotImplemented(clause);
+      llvm_unreachable(
+          "init, shutdown, set, are only valid device_num constructs");
+    }
+  }
+
+  void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, SetOp>) {
+      operation.getDefaultAsyncMutable().append(
+          createIntExpr(clause.getIntExpr()));
+    } else {
+      llvm_unreachable("set, is only valid device_num constructs");
     }
   }
 };

clang/test/CIR/CodeGenOpenACC/set.c

@@ -1,7 +1,9 @@
 // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
 
-void acc_set(void) {
-  // CHECK: cir.func @acc_set() {
+void acc_set(int cond) {
+  // CHECK: cir.func @acc_set(%[[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 set device_type(*)
   // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<star>}
@@ -10,5 +12,33 @@ void acc_set(void) {
 #pragma acc set device_type(radeon)
   // CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<radeon>}
 
+#pragma acc set default_async(cond)
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: acc.set default_async(%[[COND_CONV]] : si32)
+
+#pragma acc set default_async(1)
+  // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+  // CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32)
+
+#pragma acc set device_num(cond) if (cond)
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
+  // CHECK-NEXT: acc.set device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]])
+
+#pragma acc set device_type(radeon) default_async(1) device_num(cond) if (cond)
+  // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+  // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+  // CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+  // CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
+  // CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32) device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) attributes {device_type = #acc.device_type<radeon>}
+
   // CHECK-NEXT: cir.return
 }