[OpenACC][CIR] Add worker/vector clause lowering for Routine

clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp

@@ -323,6 +323,12 @@ class OpenACCRoutineClauseEmitter final
   void VisitSeqClause(const OpenACCSeqClause &clause) {
     routineOp.addSeq(builder.getContext(), lastDeviceTypeValues);
   }
+  void VisitWorkerClause(const OpenACCWorkerClause &clause) {
+    routineOp.addWorker(builder.getContext(), lastDeviceTypeValues);
+  }
+  void VisitVectorClause(const OpenACCVectorClause &clause) {
+    routineOp.addVector(builder.getContext(), lastDeviceTypeValues);
+  }
 };
 } // namespace
 

clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp

@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq
+void Func1() {}
+
+void Func2() {}
+#pragma acc routine(Func2) seq
+
+#pragma acc routine worker
+void Func3() {}
+
+void Func4() {}
+#pragma acc routine(Func4) worker
+
+#pragma acc routine vector
+void Func5() {}
+
+void Func6() {}
+#pragma acc routine(Func6) vector
+
+// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq
+
+// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[F3_NAME:.*Func3[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F3_R_NAME]] func(@[[F3_NAME]]) worker
+
+// CHECK: cir.func{{.*}} @[[F4_NAME:.*Func4[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F4_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[F5_NAME:.*Func5[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F5_R_NAME]] func(@[[F5_NAME]]) vector
+
+// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
+
+// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq
+// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker
+// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector

mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td

@@ -3282,6 +3282,10 @@ def OpenACC_RoutineOp : OpenACC_Op<"routine", [IsolatedFromAbove]> {
 
     // Add an entry to the 'seq' attribute for each additional device types.
     void addSeq(MLIRContext *, llvm::ArrayRef<DeviceType>);
+    // Add an entry to the 'vector' attribute for each additional device types.
+    void addVector(MLIRContext *, llvm::ArrayRef<DeviceType>);
+    // Add an entry to the 'worker' attribute for each additional device types.
+    void addWorker(MLIRContext *, llvm::ArrayRef<DeviceType>);
   }];
 
   let assemblyFormat = [{

mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp

@@ -4299,6 +4299,18 @@ void RoutineOp::addSeq(MLIRContext *context,
                                                 effectiveDeviceTypes));
 }
 
+void RoutineOp::addVector(MLIRContext *context,
+                          llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setVectorAttr(addDeviceTypeAffectedOperandHelper(context, getVectorAttr(),
+                                                   effectiveDeviceTypes));
+}
+
+void RoutineOp::addWorker(MLIRContext *context,
+                          llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setWorkerAttr(addDeviceTypeAffectedOperandHelper(context, getWorkerAttr(),
+                                                   effectiveDeviceTypes));
+}
+
 //===----------------------------------------------------------------------===//
 // InitOp
 //===----------------------------------------------------------------------===//