diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 56d4631f7845e..8e6a693841b2b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -362,6 +362,20 @@ class OpenACCRoutineClauseEmitter final
     for (const DeviceTypeArgument &arg : clause.getArchitectures())
       lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
   }
+
+  void VisitBindClause(const OpenACCBindClause &clause) {
+    if (clause.isStringArgument()) {
+      mlir::StringAttr value =
+          builder.getStringAttr(clause.getStringArgument()->getString());
+
+      routineOp.addBindStrName(builder.getContext(), lastDeviceTypeValues,
+                               value);
+    } else {
+      assert(clause.isIdentifierArgument());
+      cgm.errorNYI(clause.getSourceRange(),
+                   "Bind with an identifier argument is not yet supported");
+    }
+  }
 };
 } // namespace
 
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-bind.c b/clang/test/CIR/CodeGenOpenACC/routine-bind.c
new file mode 100644
index 0000000000000..2af024322d67e
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/routine-bind.c
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq bind("BIND1")
+void Func1(){}
+
+void Func2(){}
+#pragma acc routine(Func2) seq bind("BIND2")
+
+#pragma acc routine seq device_type(nvidia) bind("BIND3")
+void Func3(){}
+
+void Func4(){}
+#pragma acc routine(Func4) seq device_type(radeon) bind("BIND4")
+
+#pragma acc routine seq device_type(nvidia, host) bind("BIND5_N") device_type(multicore) bind("BIND5_M")
+void Func5(){}
+
+void Func6(){}
+#pragma acc routine(Func6) seq device_type(radeon) bind("BIND6_R") device_type(multicore, host) bind("BIND6_M")
+
+// 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]]) bind("BIND1") 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]]) bind("BIND3" [#acc.device_type<nvidia>]) seq
+//
+// 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]]) bind("BIND5_N" [#acc.device_type<nvidia>], "BIND5_N" [#acc.device_type<host>], "BIND5_M" [#acc.device_type<multicore>]) seq
+//
+// 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]]) bind("BIND2") seq
+// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) bind("BIND4" [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) bind("BIND6_R" [#acc.device_type<radeon>], "BIND6_M" [#acc.device_type<multicore>], "BIND6_M" [#acc.device_type<host>]) seq
+
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp b/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
new file mode 100644
index 0000000000000..2af024322d67e
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine seq bind("BIND1")
+void Func1(){}
+
+void Func2(){}
+#pragma acc routine(Func2) seq bind("BIND2")
+
+#pragma acc routine seq device_type(nvidia) bind("BIND3")
+void Func3(){}
+
+void Func4(){}
+#pragma acc routine(Func4) seq device_type(radeon) bind("BIND4")
+
+#pragma acc routine seq device_type(nvidia, host) bind("BIND5_N") device_type(multicore) bind("BIND5_M")
+void Func5(){}
+
+void Func6(){}
+#pragma acc routine(Func6) seq device_type(radeon) bind("BIND6_R") device_type(multicore, host) bind("BIND6_M")
+
+// 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]]) bind("BIND1") 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]]) bind("BIND3" [#acc.device_type<nvidia>]) seq
+//
+// 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]]) bind("BIND5_N" [#acc.device_type<nvidia>], "BIND5_N" [#acc.device_type<host>], "BIND5_M" [#acc.device_type<multicore>]) seq
+//
+// 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]]) bind("BIND2") seq
+// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) bind("BIND4" [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) bind("BIND6_R" [#acc.device_type<radeon>], "BIND6_M" [#acc.device_type<multicore>], "BIND6_M" [#acc.device_type<host>]) seq
+
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index f452686d4a30c..146dc5d087a31 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -3344,6 +3344,14 @@ def OpenACC_RoutineOp : OpenACC_Op<"routine", [IsolatedFromAbove]> {
     // Add an entry to the 'gang' attribute with a value for each additional
     // device type.
     void addGang(MLIRContext *, llvm::ArrayRef<DeviceType>, uint64_t);
+    // Add an entry to the 'bind' string-name attribute for each additional
+    // device_type.
+    void addBindStrName(MLIRContext *, llvm::ArrayRef<DeviceType>,
+                        mlir::StringAttr);
+    // Add an entry to the 'bind' ID-name attribute for each additional
+    // device_type.
+    void addBindIDName(MLIRContext *, llvm::ArrayRef<DeviceType>,
+                       mlir::SymbolRefAttr);
   }];
 
   let assemblyFormat = [{
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 64bbb1e91f293..47f122267246b 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -4464,6 +4464,45 @@ void RoutineOp::addGang(MLIRContext *context,
   setGangDimDeviceTypeAttr(mlir::ArrayAttr::get(context, deviceTypes));
 }
 
+void RoutineOp::addBindStrName(MLIRContext *context,
+                               llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+                               mlir::StringAttr val) {
+  unsigned before = getBindStrNameDeviceTypeAttr()
+                        ? getBindStrNameDeviceTypeAttr().size()
+                        : 0;
+
+  setBindStrNameDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+      context, getBindStrNameDeviceTypeAttr(), effectiveDeviceTypes));
+  unsigned after = getBindStrNameDeviceTypeAttr().size();
+
+  llvm::SmallVector<mlir::Attribute> vals;
+  if (getBindStrNameAttr())
+    llvm::copy(getBindStrNameAttr(), std::back_inserter(vals));
+  for (unsigned i = 0; i < after - before; ++i)
+    vals.push_back(val);
+
+  setBindStrNameAttr(mlir::ArrayAttr::get(context, vals));
+}
+
+void RoutineOp::addBindIDName(MLIRContext *context,
+                              llvm::ArrayRef<DeviceType> effectiveDeviceTypes,
+                              mlir::SymbolRefAttr val) {
+  unsigned before =
+      getBindIdNameDeviceTypeAttr() ? getBindIdNameDeviceTypeAttr().size() : 0;
+
+  setBindIdNameDeviceTypeAttr(addDeviceTypeAffectedOperandHelper(
+      context, getBindIdNameDeviceTypeAttr(), effectiveDeviceTypes));
+  unsigned after = getBindIdNameDeviceTypeAttr().size();
+
+  llvm::SmallVector<mlir::Attribute> vals;
+  if (getBindIdNameAttr())
+    llvm::copy(getBindIdNameAttr(), std::back_inserter(vals));
+  for (unsigned i = 0; i < after - before; ++i)
+    vals.push_back(val);
+
+  setBindIdNameAttr(mlir::ArrayAttr::get(context, vals));
+}
+
 //===----------------------------------------------------------------------===//
 // InitOp
 //===----------------------------------------------------------------------===//