[CIR][OpenACC] Implement 'routine' lowering + seq clause (#170207)

The 'routine' construct just adds a acc.routine element to the global
module, which contains all of the information about the directive. it
contains a reference to the function, which also contains a reference to
the acc.routine, which this generates.

This handles both the implicit-func version (where the routine is
    spelled without parens, and just applies to the next function) and
the explicit-func version (where the routine is spelled with the func
    name in parens).

The AST stores the directive in an OpenACCRoutineDeclAttr in the
implicit case, so we can emit that when we hit the function declaration.
The explicit case is held in an OpenACCRoutineAnnotAttr on the function,
however, when we emit the function we haven't necessarily seen the
construct yet, so we can't depend on that attribute. Instead, we save up
the list in Sema so that we can emit them all at the end.

This results in the tests getting really hard to read (because ordering
is a little awkward based on spelling, with no way to fix it), so we
instead split the tests up based on topic.

One last thing: Flang spends some time determining if the clause lists
of two routines on the same function are identical, and omits the
duplicates. However, it seems to do a poor job on this when the ordering
isn't the same, or references are slightly different. This patch doesn't
bother trying that, and instead emits all, trusting the ACC dialect to
remove duplicates/handle duplicates gracefully.

Note; This doesn't cause emission of functions that would otherwise not
be emitted, but DOES emit routine references based on which function
they are attached to.

clang/include/clang/AST/ASTConsumer.h

@@ -27,6 +27,7 @@ namespace clang {
   class VarDecl;
   class FunctionDecl;
   class ImportDecl;
+  class OpenACCRoutineDecl;
 
 /// ASTConsumer - This is an abstract interface that should be implemented by
 /// clients that read ASTs.  This abstraction layer allows the client to be
@@ -116,6 +117,11 @@ public:
   // variable has been instantiated.
   virtual void HandleCXXStaticMemberVarInstantiation(VarDecl *D) {}
 
+  /// Callback to handle the end-of-translation unit attachment of OpenACC
+  /// routine declaration information.
+  virtual void HandleOpenACCRoutineReference(const FunctionDecl *FD,
+                                             const OpenACCRoutineDecl *RD) {}
+
   /// Callback involved at the end of a translation unit to
   /// notify the consumer that a vtable for the given C++ class is
   /// required.

clang/include/clang/CIR/CIRGenerator.h

@@ -81,6 +81,9 @@ public:
   void HandleTagDeclDefinition(clang::TagDecl *d) override;
   void HandleTagDeclRequiredDefinition(const clang::TagDecl *D) override;
   void HandleCXXStaticMemberVarInstantiation(clang::VarDecl *D) override;
+  void
+  HandleOpenACCRoutineReference(const clang::FunctionDecl *FD,
+                                const clang::OpenACCRoutineDecl *RD) override;
   void CompleteTentativeDefinition(clang::VarDecl *d) override;
   void HandleVTable(clang::CXXRecordDecl *rd) override;
 

clang/include/clang/Sema/SemaOpenACC.h

@@ -37,8 +37,16 @@ class Scope;
 class SemaOpenACC : public SemaBase {
 public:
   using DeclGroupPtrTy = OpaquePtr<DeclGroupRef>;
+  using RoutineRefListTy = std::pair<FunctionDecl *, OpenACCRoutineDecl *>;
 
 private:
+  // We save a list of routine clauses that refer to a different function(that
+  // is, routine-with-a-name) so that we can do the emission at the 'end'.  We
+  // have to do this, since functions can be emitted before they are referenced,
+  // and the OpenACCRoutineDecl isn't necessarily emitted, as it might be in a
+  // function/etc. So we do these emits at the end of the TU.
+  llvm::SmallVector<RoutineRefListTy> RoutineRefList;
+
   struct ComputeConstructInfo {
     /// Which type of compute construct we are inside of, which we can use to
     /// determine whether we should add loops to the above collection.  We can
@@ -752,6 +760,7 @@ public:
   };
 
   SemaOpenACC(Sema &S);
+  void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU);
 
   // Called when we encounter a 'while' statement, before looking at its 'body'.
   void ActOnWhileStmt(SourceLocation WhileLoc);

clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp

@@ -287,9 +287,82 @@ void CIRGenModule::emitGlobalOpenACCDeclareDecl(const OpenACCDeclareDecl *d) {
 }
 
 void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
-  getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Routine Construct");
+  // Do nothing here. The OpenACCRoutineDeclAttr handles the implicit name
+  // cases, and the end-of-TU handling manages the named cases. This is
+  // necessary because these references aren't necessarily emitted themselves,
+  // but can be named anywhere.
 }
 
 void CIRGenModule::emitGlobalOpenACCRoutineDecl(const OpenACCRoutineDecl *d) {
-  errorNYI(d->getSourceRange(), "OpenACC Global Routine Construct");
+  // Do nothing here. The OpenACCRoutineDeclAttr handles the implicit name
+  // cases, and the end-of-TU handling manages the named cases. This is
+  // necessary because these references aren't necessarily emitted themselves,
+  // but can be named anywhere.
+}
+
+namespace {
+class OpenACCRoutineClauseEmitter final
+    : public OpenACCClauseVisitor<OpenACCRoutineClauseEmitter> {
+  CIRGen::CIRGenBuilderTy &builder;
+  mlir::acc::RoutineOp routineOp;
+  llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
+
+public:
+  OpenACCRoutineClauseEmitter(CIRGen::CIRGenBuilderTy &builder,
+                              mlir::acc::RoutineOp routineOp)
+      : builder(builder), routineOp(routineOp) {}
+
+  void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
+    this->VisitClauseList(clauses);
+  }
+
+  void VisitClause(const OpenACCClause &clause) {
+    llvm_unreachable("Invalid OpenACC clause on routine");
+  }
+
+  void VisitSeqClause(const OpenACCSeqClause &clause) {
+    routineOp.addSeq(builder.getContext(), lastDeviceTypeValues);
+  }
+};
+} // namespace
+
+void CIRGenModule::emitOpenACCRoutineDecl(
+    const clang::FunctionDecl *funcDecl, cir::FuncOp func,
+    SourceLocation pragmaLoc, ArrayRef<const OpenACCClause *> clauses) {
+  mlir::OpBuilder::InsertionGuard guardCase(builder);
+  // These need to appear at the global module.
+  builder.setInsertionPointToEnd(&getModule().getBodyRegion().front());
+
+  mlir::Location routineLoc = getLoc(pragmaLoc);
+
+  std::stringstream routineNameSS;
+  // This follows the same naming format as Flang.
+  routineNameSS << "acc_routine_" << routineCounter++;
+  std::string routineName = routineNameSS.str();
+
+  // There isn't a good constructor for RoutineOp that just takes a location +
+  // name + function, so we use one that creates an otherwise RoutineOp and
+  // count on the visitor/emitter to fill these in.
+  auto routineOp = mlir::acc::RoutineOp::create(
+      builder, routineLoc, routineName,
+      mlir::SymbolRefAttr::get(builder.getContext(), func.getName()),
+      /*implicit=*/false);
+
+  // We have to add a pointer going the other direction via an acc.routine_info,
+  // from the func to the routine.
+  llvm::SmallVector<mlir::SymbolRefAttr> funcRoutines;
+  if (auto routineInfo =
+          func.getOperation()->getAttrOfType<mlir::acc::RoutineInfoAttr>(
+              mlir::acc::getRoutineInfoAttrName()))
+    funcRoutines.append(routineInfo.getAccRoutines().begin(),
+                        routineInfo.getAccRoutines().end());
+
+  funcRoutines.push_back(
+      mlir::SymbolRefAttr::get(builder.getContext(), routineName));
+  func.getOperation()->setAttr(
+      mlir::acc::getRoutineInfoAttrName(),
+      mlir::acc::RoutineInfoAttr::get(func.getContext(), funcRoutines));
+
+  OpenACCRoutineClauseEmitter emitter{builder, routineOp};
+  emitter.emitClauses(clauses);
 }

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -2227,6 +2227,15 @@ CIRGenModule::createCIRFunction(mlir::Location loc, StringRef name,
 
     if (!cgf)
       theModule.push_back(func);
+
+    if (this->getLangOpts().OpenACC) {
+      // We only have to handle this attribute, since OpenACCAnnotAttrs are
+      // handled via the end-of-TU work.
+      for (const auto *attr :
+           funcDecl->specific_attrs<OpenACCRoutineDeclAttr>())
+        emitOpenACCRoutineDecl(funcDecl, func, attr->getLocation(),
+                               attr->Clauses);
+    }
   }
   return func;
 }

clang/lib/CIR/CodeGen/CIRGenModule.h

@@ -461,6 +461,12 @@ public:
                                             OpenACCModifierKind modifiers,
                                             bool structured, bool implicit,
                                             bool requiresDtor);
+  // Each of the acc.routine operations must have a unique name, so we just use
+  // an integer counter.  This is how Flang does it, so it seems reasonable.
+  unsigned routineCounter = 0;
+  void emitOpenACCRoutineDecl(const clang::FunctionDecl *funcDecl,
+                              cir::FuncOp func, SourceLocation pragmaLoc,
+                              ArrayRef<const OpenACCClause *> clauses);
 
   // C++ related functions.
   void emitDeclContext(const DeclContext *dc);

clang/lib/CIR/CodeGen/CIRGenerator.cpp

@@ -166,6 +166,18 @@ void CIRGenerator::HandleCXXStaticMemberVarInstantiation(VarDecl *D) {
   cgm->handleCXXStaticMemberVarInstantiation(D);
 }
 
+void CIRGenerator::HandleOpenACCRoutineReference(const FunctionDecl *FD,
+                                                 const OpenACCRoutineDecl *RD) {
+  llvm::StringRef mangledName = cgm->getMangledName(FD);
+  cir::FuncOp entry =
+      mlir::dyn_cast_if_present<cir::FuncOp>(cgm->getGlobalValue(mangledName));
+
+  // if this wasn't generated, don't force it to be.
+  if (!entry)
+    return;
+  cgm->emitOpenACCRoutineDecl(FD, entry, RD->getBeginLoc(), RD->clauses());
+}
+
 void CIRGenerator::CompleteTentativeDefinition(VarDecl *d) {
   if (diags.hasErrorOccurred())
     return;

clang/lib/CIR/FrontendAction/CIRGenAction.cpp

@@ -88,6 +88,11 @@ public:
     Gen->HandleCXXStaticMemberVarInstantiation(VD);
   }
 
+  void HandleOpenACCRoutineReference(const FunctionDecl *FD,
+                                     const OpenACCRoutineDecl *RD) override {
+    Gen->HandleOpenACCRoutineReference(FD, RD);
+  }
+
   void HandleInlineFunctionDefinition(FunctionDecl *D) override {
     Gen->HandleInlineFunctionDefinition(D);
   }

clang/lib/Sema/Sema.cpp

@@ -1497,6 +1497,9 @@ void Sema::ActOnEndOfTranslationUnit() {
 
   if (LangOpts.HLSL)
     HLSL().ActOnEndOfTranslationUnit(getASTContext().getTranslationUnitDecl());
+  if (LangOpts.OpenACC)
+    OpenACC().ActOnEndOfTranslationUnit(
+        getASTContext().getTranslationUnitDecl());
 
   // If there were errors, disable 'unused' warnings since they will mostly be
   // noise. Don't warn for a use from a module: either we should warn on all

clang/lib/Sema/SemaOpenACC.cpp

@@ -12,6 +12,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "clang/Sema/SemaOpenACC.h"
+#include "clang/AST/ASTConsumer.h"
 #include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/StmtOpenACC.h"
 #include "clang/Basic/DiagnosticSema.h"
@@ -2457,7 +2458,8 @@ OpenACCRoutineDecl *SemaOpenACC::CheckRoutineDecl(
     ArrayRef<const OpenACCClause *> Clauses, SourceLocation EndLoc) {
   assert(LParenLoc.isValid());
 
-  if (FunctionDecl *FD = getFunctionFromRoutineName(FuncRef)) {
+  FunctionDecl *FD = nullptr;
+  if ((FD = getFunctionFromRoutineName(FuncRef))) {
     // OpenACC 3.3 2.15:
     // In C and C++, function static variables are not supported in functions to
     // which a routine directive applies.
@@ -2509,11 +2511,9 @@ OpenACCRoutineDecl *SemaOpenACC::CheckRoutineDecl(
                                                         {DirLoc, BindLoc});
     FD->addAttr(RAA);
     // In case we are referencing not the 'latest' version, make sure we add
-    // the attribute to all declarations.
-    while (FD != FD->getMostRecentDecl()) {
-      FD = FD->getMostRecentDecl();
-      FD->addAttr(RAA);
-    }
+    // the attribute to all declarations after the 'found' one.
+    for (auto *CurFD : FD->redecls())
+      CurFD->addAttr(RAA->clone(getASTContext()));
   }
 
   LastRoutineDecl = OpenACCRoutineDecl::Create(
@@ -2522,9 +2522,20 @@ OpenACCRoutineDecl *SemaOpenACC::CheckRoutineDecl(
   LastRoutineDecl->setAccess(AS_public);
   getCurContext()->addDecl(LastRoutineDecl);
 
+  if (FD) {
+    // Add this attribute to the list of annotations so that codegen can visit
+    // it later. FD doesn't necessarily exist, but that case should be
+    // diagnosed.
+    RoutineRefList.emplace_back(FD, LastRoutineDecl);
+  }
   return LastRoutineDecl;
 }
 
+void SemaOpenACC::ActOnEndOfTranslationUnit(TranslationUnitDecl *TU) {
+  for (auto [FD, RoutineDecl] : RoutineRefList)
+    SemaRef.Consumer.HandleOpenACCRoutineReference(FD, RoutineDecl);
+}
+
 DeclGroupRef SemaOpenACC::ActOnEndRoutineDeclDirective(
     SourceLocation StartLoc, SourceLocation DirLoc, SourceLocation LParenLoc,
     Expr *ReferencedFunc, SourceLocation RParenLoc,

clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp

@@ -1,6 +0,0 @@
-// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
-// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-llvm %s -o %t-cir.ll -verify
-
-void foo() {}
-// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Global Routine Construct}}
-#pragma acc routine(foo) seq

clang/test/CIR/CodeGenOpenACC/routine-anon-ns.cpp

@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+namespace {
+#pragma acc routine seq
+  void NSFunc1(){}
+#pragma acc routine seq
+  auto Lambda1 = [](){};
+
+  auto Lambda2 = [](){};
+} // namespace 
+
+#pragma acc routine(NSFunc1) seq
+#pragma acc routine(Lambda2) seq
+void force_emit() {
+  NSFunc1();
+  Lambda1();
+  Lambda2();
+}
+
+// CHECK: cir.func{{.*}} @[[F1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]], @[[F1_R2_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L2_R_NAME:.*]]]>}
+//
+// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq
+// CHECK: acc.routine @[[L1_R_NAME]] func(@[[L1_NAME]]) seq
+// CHECK: acc.routine @[[F1_R2_NAME]] func(@[[F1_NAME]]) seq
+// CHECK: acc.routine @[[L2_R_NAME]] func(@[[L2_NAME]]) seq

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

@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq
+auto Lambda1 = [](){};
+
+auto Lambda2 = [](){};
+#pragma acc routine(Lambda2) seq
+#pragma acc routine(Lambda2) seq
+
+#pragma acc routine seq
+int GlobalFunc1();
+
+int GlobalFunc2();
+#pragma acc routine(GlobalFunc2) seq
+#pragma acc routine(GlobalFunc1) seq
+
+void force_emit() {
+  Lambda1();
+  Lambda2();
+  GlobalFunc1();
+  GlobalFunc2();
+}
+
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L2_R_NAME:.*]], @[[L2_R2_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[G1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G1_R_NAME:.*]], @[[G1_R2_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[G2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G2_R_NAME:.*]]]>}
+
+// CHECK: acc.routine @[[L1_R_NAME]] func(@[[L1_NAME]]) seq
+// CHECK: acc.routine @[[G1_R_NAME]] func(@[[G1_NAME]]) seq
+// CHECK: acc.routine @[[L2_R_NAME]] func(@[[L2_NAME]]) seq
+// CHECK: acc.routine @[[L2_R2_NAME]] func(@[[L2_NAME]]) seq
+// CHECK: acc.routine @[[G2_R_NAME]] func(@[[G2_NAME]]) seq
+// CHECK: acc.routine @[[G1_R2_NAME]] func(@[[G1_NAME]]) seq

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

@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq
+void GlobalFunc4();
+#pragma acc routine(GlobalFunc4) seq
+
+#pragma acc routine seq
+#pragma acc routine seq
+void GlobalFunc5();
+#pragma acc routine(GlobalFunc5) seq
+#pragma acc routine(GlobalFunc5) seq
+
+void GlobalFunc6();
+void GlobalFunc6();
+#pragma acc routine(GlobalFunc6) seq
+void GlobalFunc6(){}
+
+void GlobalFunc7(){}
+#pragma acc routine(GlobalFunc7) seq
+
+void force_emit() {
+  GlobalFunc4();
+  GlobalFunc5();
+  GlobalFunc6();
+  GlobalFunc7();
+}
+
+// CHECK: cir.func{{.*}} @[[G6_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G6_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[G7_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G7_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[G4_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G4_R_NAME:.*]], @[[G4_R2_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[G5_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G5_R_NAME:.*]], @[[G5_R1_NAME:.*]], @[[G5_R2_NAME:.*]], @[[G5_R3_NAME:.*]]]>}
+
+// CHECK: acc.routine @[[G4_R_NAME]] func(@[[G4_NAME]]) seq
+// CHECK: acc.routine @[[G5_R_NAME]] func(@[[G5_NAME]]) seq
+// CHECK: acc.routine @[[G5_R1_NAME]] func(@[[G5_NAME]]) seq
+//
+// CHECK: acc.routine @[[G4_R2_NAME]] func(@[[G4_NAME]]) seq
+//
+// CHECK: acc.routine @[[G5_R2_NAME]] func(@[[G5_NAME]]) seq
+// CHECK: acc.routine @[[G5_R3_NAME]] func(@[[G5_NAME]]) seq
+//
+// CHECK: acc.routine @[[G6_R_NAME]] func(@[[G6_NAME]]) seq
+// CHECK: acc.routine @[[G7_R_NAME]] func(@[[G7_NAME]]) seq

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

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+void GlobalFunc();
+void InFunc() {
+
+#pragma acc routine(GlobalFunc) seq
+  GlobalFunc();
+
+#pragma acc routine seq
+  auto Lambda1 = [](){};
+  Lambda1();
+
+  auto Lambda2 = [](){};
+#pragma acc routine(Lambda2) seq
+  Lambda2();
+};
+
+// CHECK: cir.func{{.*}} @[[G1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[G1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L2_R_NAME:.*]]]>}
+
+// CHECK: acc.routine @[[L1_R_NAME]] func(@[[L1_NAME]]) seq
+// CHECK: acc.routine @[[G1_R_NAME]] func(@[[G1_NAME]]) seq
+// CHECK: acc.routine @[[L2_R_NAME]] func(@[[L2_NAME]]) seq

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

@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+struct S {
+#pragma acc routine seq
+  void MemFunc1();
+  void MemFunc2();
+#pragma acc routine(S::MemFunc2) seq
+  void MemFunc3();
+#pragma acc routine(S::MemFunc3) seq
+
+#pragma acc routine seq
+  static void StaticMemFunc1();
+  static void StaticMemFunc2();
+  static void StaticMemFunc3();
+#pragma acc routine(StaticMemFunc3) seq
+
+#pragma acc routine seq
+  static constexpr auto StaticLambda1 = [](){};
+ static constexpr auto StaticLambda2 = [](){};
+};
+#pragma acc routine(S::MemFunc2) seq
+#pragma acc routine(S::StaticLambda2) seq
+#pragma acc routine(S::StaticMemFunc2) seq
+
+void force_emit() {
+  S{}.MemFunc1();
+  S{}.MemFunc2();
+  S{}.MemFunc3();
+  S::StaticMemFunc1();
+  S::StaticMemFunc2();
+  S::StaticMemFunc3();
+  S::StaticLambda1();
+  S::StaticLambda2();
+}
+
+// CHECK: cir.func{{.*}} @[[MEM1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[MEM1_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[MEM2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[MEM2_R_NAME:.*]], @[[MEM2_R2_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[MEM3_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[MEM3_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[STATICMEM1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[STATICMEM1_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[STATICMEM2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[STATICMEM2_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[STATICMEM3_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[STATICMEM3_R_NAME:.*]]]>}
+//
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L2_R_NAME:.*]]]>}
+//
+// CHECK: acc.routine @[[MEM1_R_NAME]] func(@[[MEM1_NAME]]) seq
+// CHECK: acc.routine @[[STATICMEM1_R_NAME]] func(@[[STATICMEM1_NAME]]) seq
+// CHECK: acc.routine @[[L1_R_NAME]] func(@[[L1_NAME]]) seq
+// CHECK: acc.routine @[[MEM2_R_NAME]] func(@[[MEM2_NAME]]) seq
+// CHECK: acc.routine @[[MEM3_R_NAME]] func(@[[MEM3_NAME]]) seq
+// CHECK: acc.routine @[[STATICMEM3_R_NAME]] func(@[[STATICMEM3_NAME]]) seq
+// CHECK: acc.routine @[[MEM2_R2_NAME]] func(@[[MEM2_NAME]]) seq
+// CHECK: acc.routine @[[L2_R_NAME]] func(@[[L2_NAME]]) seq
+// CHECK: acc.routine @[[STATICMEM2_R_NAME]] func(@[[STATICMEM2_NAME]]) seq

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

@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+namespace NS1 {
+#pragma acc routine seq
+  int NSFunc1();
+#pragma acc routine seq
+  auto Lambda1 = [](){};
+
+  auto Lambda2 = [](){};
+} // namespace NS1
+
+#pragma acc routine(NS1::NSFunc1) seq
+#pragma acc routine(NS1::Lambda2) seq
+
+void force_emit() {
+  NS1::NSFunc1();
+  NS1::Lambda1();
+  NS1::Lambda2();
+}
+
+// CHECK: cir.func{{.*}} @[[F1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]], @[[F1_R2_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L1_R_NAME:.*]]]>}
+// CHECK: cir.func {{.*}}lambda{{.*}} @[[L2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[L2_R_NAME:.*]]]>}
+//
+// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq 
+// CHECK: acc.routine @[[L1_R_NAME]] func(@[[L1_NAME]]) seq 
+// CHECK: acc.routine @[[F1_R2_NAME]] func(@[[F1_NAME]]) seq  
+// CHECK: acc.routine @[[L2_R_NAME]] func(@[[L2_NAME]]) seq 

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

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq
+template<typename T>
+void func(){}
+
+void use() {
+  func<int>();
+  func<float>();
+}
+
+// CHECK: cir.func{{.*}} @[[T1_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[T1_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @[[T2_NAME:[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[T2_R_NAME:.*]]]>}
+//
+// CHECK: acc.routine @[[T1_R_NAME]] func(@[[T1_NAME]]) seq
+// CHECK: acc.routine @[[T2_R_NAME]] func(@[[T2_NAME]]) seq

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

@@ -3232,6 +3232,18 @@ def OpenACC_RoutineOp : OpenACC_Op<"routine", [IsolatedFromAbove]> {
       OptionalAttr<DeviceTypeArrayAttr>:$gangDimDeviceType);
 
   let extraClassDeclaration = [{
+    // 'create' function to generate an 'empty' routine.
+    static RoutineOp create(::mlir::OpBuilder & builder,
+                            ::mlir::Location location,
+                            ::llvm::StringRef sym_name,
+                            mlir::SymbolRefAttr func_name, bool implicit) {
+      return create(builder, location, sym_name, func_name, /*bindIDName=*/{},
+                    /*bindStrName=*/{}, /*bindIdNameDeviceType=*/{},
+                    /*bindStrnameDeviceType=*/{}, /*worker=*/{}, /*vector=*/{},
+                    /*seq=*/{}, /*nohost=*/false, implicit, /*gang=*/{},
+                    /*gangDim=*/{}, /*gangDimDeviceType=*/{});
+    }
+
     static StringRef getGangDimKeyword() { return "dim"; }
 
     /// Return true if the op has the worker attribute for the
@@ -3267,6 +3279,9 @@ def OpenACC_RoutineOp : OpenACC_Op<"routine", [IsolatedFromAbove]> {
 
     std::optional<::std::variant<mlir::SymbolRefAttr, mlir::StringAttr>> getBindNameValue();
     std::optional<::std::variant<mlir::SymbolRefAttr, mlir::StringAttr>> getBindNameValue(mlir::acc::DeviceType deviceType);
+
+    // Add an entry to the 'seq' attribute for each additional device types.
+    void addSeq(MLIRContext *, llvm::ArrayRef<DeviceType>);
   }];
 
   let assemblyFormat = [{

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

@@ -4293,6 +4293,12 @@ RoutineOp::getGangDimValue(mlir::acc::DeviceType deviceType) {
   return std::nullopt;
 }
 
+void RoutineOp::addSeq(MLIRContext *context,
+                       llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+  setSeqAttr(addDeviceTypeAffectedOperandHelper(context, getSeqAttr(),
+                                                effectiveDeviceTypes));
+}
+
 //===----------------------------------------------------------------------===//
 // InitOp
 //===----------------------------------------------------------------------===//