Author: Erich Keane Date: 2025-12-08T06:23:13-08:00 New Revision: dd06214394977729a3f2715bfadf1b31467551b8
URL: https://github.com/llvm/llvm-project/commit/dd06214394977729a3f2715bfadf1b31467551b8 DIFF: https://github.com/llvm/llvm-project/commit/dd06214394977729a3f2715bfadf1b31467551b8.diff LOG: [OpenACC][CIR] Implement routine 'bind'-with-a-string lowering (#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. Added: clang/test/CIR/CodeGenOpenACC/routine-bind.c clang/test/CIR/CodeGenOpenACC/routine-bind.cpp Modified: clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp Removed: ################################################################################ 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
