[OpenACC][CIR] Implement Loop lowering of seq/auto/independent

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.h

@@ -147,13 +147,13 @@ class OpenACCClauseCIREmitter final
             decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
     } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
                                       mlir::acc::SerialOp, mlir::acc::KernelsOp,
-                                      mlir::acc::DataOp>) {
+                                      mlir::acc::DataOp, mlir::acc::LoopOp>) {
       // Nothing to do here, these constructs don't have any IR for these, as
       // they just modify the other clauses IR.  So setting of
       // `lastDeviceTypeValues` (done above) is all we need.
     } else {
       // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. update, data, loop, routine, combined constructs remain.
+      // unreachable. update, data, routine, combined constructs remain.
       return clauseNotImplemented(clause);
     }
   }
@@ -306,6 +306,36 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("set, is only valid device_num constructs");
     }
   }
+
+  void VisitSeqClause(const OpenACCSeqClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+      operation.addSeq(builder.getContext(), lastDeviceTypeValues);
+    } else {
+      // TODO: When we've implemented this for everything, switch this to an
+      // unreachable. Routine, Combined constructs remain.
+      return clauseNotImplemented(clause);
+    }
+  }
+
+  void VisitAutoClause(const OpenACCAutoClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+      operation.addAuto(builder.getContext(), lastDeviceTypeValues);
+    } else {
+      // TODO: When we've implemented this for everything, switch this to an
+      // unreachable. Routine, Combined constructs remain.
+      return clauseNotImplemented(clause);
+    }
+  }
+
+  void VisitIndependentClause(const OpenACCIndependentClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
+      operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
+    } else {
+      // TODO: When we've implemented this for everything, switch this to an
+      // unreachable. Routine, Combined constructs remain.
+      return clauseNotImplemented(clause);
+    }
+  }
 };
 
 template <typename OpTy>

clang/test/CIR/CodeGenOpenACC/loop.cpp

@@ -30,4 +30,83 @@ extern "C" void acc_loop(int *A, int *B, int *C, int N) {
   // CHECK-NEXT: } loc
   // CHECK-NEXT: acc.yield
   // CHECK-NEXT: } loc
+
+
+#pragma acc loop seq
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) seq
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) seq
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<radeon>]} loc
+#pragma acc loop seq device_type(nvidia, radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+#pragma acc loop seq device_type(radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc
+
+#pragma acc loop independent
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) independent
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {independent = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) independent
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {independent = [#acc.device_type<radeon>]} loc
+#pragma acc loop independent device_type(nvidia, radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+#pragma acc loop independent device_type(radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc
+
+#pragma acc loop auto
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
+#pragma acc loop device_type(nvidia, radeon) auto
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc
+#pragma acc loop device_type(radeon) auto
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<radeon>]} loc
+#pragma acc loop auto device_type(nvidia, radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
+#pragma acc loop auto device_type(radeon)
+  for(unsigned I = 0; I < N; ++I);
+  // CHECK: acc.loop {
+  // CHECK: acc.yield
+  // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc
 }

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

@@ -2198,6 +2198,14 @@ def OpenACC_LoopOp : OpenACC_Op<"loop",
     /// Return the value of the worker clause for the given device_type 
     /// if present.
     mlir::Value getGangValue(mlir::acc::GangArgType gangArgType, mlir::acc::DeviceType deviceType);
+
+    // Add an entry to the 'seq' attribute for each additional device types.
+    void addSeq(MLIRContext *, llvm::ArrayRef<DeviceType>);
+    // Add an entry to the 'independent' attribute for each additional device
+    // types.
+    void addIndependent(MLIRContext *, llvm::ArrayRef<DeviceType>);
+    // Add an entry to the 'auto' attribute for each additional device types.
+    void addAuto(MLIRContext *, llvm::ArrayRef<DeviceType>);
   }];
 
   let hasCustomAssemblyFormat = 1;

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

@@ -2651,6 +2651,24 @@ void printLoopControl(OpAsmPrinter &p, Operation *op, Region &region,
   p.printRegion(region, /*printEntryBlockArgs=*/false);
 }
 
+void acc::LoopOp::addSeq(MLIRContext *context,
+                         llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setSeqAttr(addDeviceTypeAffectedOperandHelper(context, getSeqAttr(),
+                                                effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addIndependent(
+    MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setIndependentAttr(addDeviceTypeAffectedOperandHelper(
+      context, getIndependentAttr(), effectiveDeviceTypes));
+}
+
+void acc::LoopOp::addAuto(MLIRContext *context,
+                          llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setAuto_Attr(addDeviceTypeAffectedOperandHelper(context, getAuto_Attr(),
+                                                  effectiveDeviceTypes));
+}
+
 //===----------------------------------------------------------------------===//
 // DataOp
 //===----------------------------------------------------------------------===//