https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/170916
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.
>From 2a2f42be6481eefa1f293f2e0db21c19265c4049 Mon Sep 17 00:00:00 2001
From: erichkeane <[email protected]>
Date: Fri, 5 Dec 2025 11:38:52 -0800
Subject: [PATCH] [OpenACC][CIR] Implement routine 'bind'-with-a-string
lowering
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.
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 14 +++++++
clang/test/CIR/CodeGenOpenACC/routine-bind.c | 39 +++++++++++++++++++
.../test/CIR/CodeGenOpenACC/routine-bind.cpp | 39 +++++++++++++++++++
.../mlir/Dialect/OpenACC/OpenACCOps.td | 8 ++++
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 39 +++++++++++++++++++
5 files changed, 139 insertions(+)
create mode 100644 clang/test/CIR/CodeGenOpenACC/routine-bind.c
create mode 100644 clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
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
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
