llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
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.
---
Full diff: https://github.com/llvm/llvm-project/pull/170916.diff
5 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp (+14)
- (added) clang/test/CIR/CodeGenOpenACC/routine-bind.c (+39)
- (added) clang/test/CIR/CodeGenOpenACC/routine-bind.cpp (+39)
- (modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+8)
- (modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+39)
``````````diff
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
//===----------------------------------------------------------------------===//
``````````
</details>
https://github.com/llvm/llvm-project/pull/170916
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
