[OpenACC][CIR] Implement routine 'bind'-with-a-string lowering

erichkeane · erichkeane · commit 2a2f42be6481 · 2025-12-05T11:41:22.000-08:00
The 'bind' clause emits an attribute on the RoutineOp that states which
function it should call on the device side.  When provided in double-quotes, the
function on the device side should be the exact name given. This patch
emits the IR to do that.

As a part of that, we add a helper function to the OpenACC dialect to do
so, as well as a version that adds the ID version (though we don't
    exercise th at yet).

The 'bind' with an ID should do the MANGLED name, but it isn't quite
clear what that name SHOULD be yet.  Since the signature of a function
is included in its mangling, and we're not providing said signature, we
have to come up with something.  This is left as an exercise for a
future patch.
diff --git 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
@@ -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
@@ -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
@@ -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
@@ -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
 //===----------------------------------------------------------------------===//
