[OpenACC][CIR] 'bind' lowering with identifier #171749
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
The bind clause specifies the name of the function to call on the device, and takes either a string or identifier(per the standard): "If the name is specified as an identifier, it is callled as if the name were specified in the language being compiled. If the name is specified as a string, the string is used for the procedure name unmodified". The latter (as a string) is already implemented, this patch implements the former. Unfortunately, no existing implementation of this in C++ seems to exist. Other languages, the 'name' of a function is sufficient to identify it (in this case 'bind' can refer to undeclared functions), so it is possible to figure out what the name should be. In C++ with overloading (without a discriminator, ala-fortran), a name only names an infinite overload set. SO, in order to implement this, I've decided that the 'called as' (bound) function must have the same signature as the one marked by the 'routine'. This is trivially sensible in non-member functions, however requires a bit more thought for member(and thus lambda-call-operators) functions. In this case, we 'promote' the type of the function to a 'free' function by turning the implicit 'this' to an explicit 'this'. I believe this is the most sensible and reasonable way to implement this, and really the only way to make something usable.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThe bind clause specifies the name of the function to call on the device, and takes either a string or identifier(per the standard): "If the name is specified as an identifier, it is callled as if the name were specified in the language being compiled. If the name is specified as a string, the string is used for the procedure name unmodified". The latter (as a string) is already implemented, this patch implements the former. Unfortunately, no existing implementation of this in C++ seems to exist. Other languages, the 'name' of a function is sufficient to identify it (in this case 'bind' can refer to undeclared functions), so it is possible to figure out what the name should be. In C++ with overloading (without a discriminator, ala-fortran), a name only names an infinite overload set. SO, in order to implement this, I've decided that the 'called as' (bound) function must have the same signature as the one marked by the 'routine'. This is trivially sensible in non-member functions, however requires a bit more thought for member(and thus lambda-call-operators) functions. In this case, we 'promote' the type of the function to a 'free' function by turning the implicit 'this' to an explicit 'this'. I believe this is the most sensible and reasonable way to implement this, and really the only way to make something usable. Full diff: https://github.com/llvm/llvm-project/pull/171749.diff 5 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 8e6a693841b2b..87b6596eb6773 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -306,13 +306,15 @@ class OpenACCRoutineClauseEmitter final
CIRGenModule &cgm;
CIRGen::CIRGenBuilderTy &builder;
mlir::acc::RoutineOp routineOp;
+ const clang::FunctionDecl *funcDecl;
llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
public:
OpenACCRoutineClauseEmitter(CIRGenModule &cgm,
CIRGen::CIRGenBuilderTy &builder,
- mlir::acc::RoutineOp routineOp)
- : cgm(cgm), builder(builder), routineOp(routineOp) {}
+ mlir::acc::RoutineOp routineOp,
+ const clang::FunctionDecl *funcDecl)
+ : cgm(cgm), builder(builder), routineOp(routineOp), funcDecl(funcDecl) {}
void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
this->VisitClauseList(clauses);
@@ -372,8 +374,12 @@ class OpenACCRoutineClauseEmitter final
value);
} else {
assert(clause.isIdentifierArgument());
- cgm.errorNYI(clause.getSourceRange(),
- "Bind with an identifier argument is not yet supported");
+ std::string bindName = cgm.getOpenACCBindMangledName(
+ clause.getIdentifierArgument(), funcDecl);
+
+ routineOp.addBindIDName(
+ builder.getContext(), lastDeviceTypeValues,
+ mlir::SymbolRefAttr::get(builder.getContext(), bindName));
}
}
};
@@ -416,6 +422,6 @@ void CIRGenModule::emitOpenACCRoutineDecl(
mlir::acc::getRoutineInfoAttrName(),
mlir::acc::RoutineInfoAttr::get(func.getContext(), funcRoutines));
- OpenACCRoutineClauseEmitter emitter{*this, builder, routineOp};
+ OpenACCRoutineClauseEmitter emitter{*this, builder, routineOp, funcDecl};
emitter.emitClauses(clauses);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 1ad1c2fa41aa1..ed8efd64fe8f1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1721,6 +1721,70 @@ static std::string getMangledNameImpl(CIRGenModule &cgm, GlobalDecl gd,
return std::string(out.str());
}
+static FunctionDecl *
+createOpenACCBindTempFunction(ASTContext &ctx, const IdentifierInfo *bindName,
+ const FunctionDecl *protoFunc) {
+ // If this is a C no-prototype function, we can take the 'easy' way out and
+ // just create a function with no arguments/functions, etc.
+ if (!protoFunc->hasPrototype())
+ return FunctionDecl::Create(
+ ctx, /*DC=*/ctx.getTranslationUnitDecl(),
+ /*StartLoc=*/SourceLocation{}, /*NLoc=*/SourceLocation{}, bindName,
+ protoFunc->getType(), /*TInfo=*/nullptr, StorageClass::SC_None);
+
+ QualType funcTy = protoFunc->getType();
+ auto *FPT = cast<FunctionProtoType>(protoFunc->getType());
+
+ // If this is a member function, add an explicit 'this' to the function type.
+ if (auto *methodDecl = dyn_cast<CXXMethodDecl>(protoFunc);
+ methodDecl && methodDecl->isImplicitObjectMemberFunction()) {
+ llvm::SmallVector<QualType> paramTypes{FPT->getParamTypes()};
+ paramTypes.insert(paramTypes.begin(), methodDecl->getThisType());
+
+ funcTy = ctx.getFunctionType(FPT->getReturnType(), paramTypes,
+ FPT->getExtProtoInfo());
+ FPT = cast<FunctionProtoType>(funcTy);
+ }
+
+ auto *tempFunc =
+ FunctionDecl::Create(ctx, /*DC=*/ctx.getTranslationUnitDecl(),
+ /*StartLoc=*/SourceLocation{},
+ /*NLoc=*/SourceLocation{}, bindName, funcTy,
+ /*TInfo=*/nullptr, StorageClass::SC_None);
+
+ SmallVector<ParmVarDecl *, 16> params;
+
+ // Add all of the parameters.
+ for (unsigned i = 0, e = FPT->getNumParams(); i != e; ++i) {
+ ParmVarDecl *parm = ParmVarDecl::Create(
+ ctx, tempFunc, /*StartLoc=*/SourceLocation{},
+ /*IdLoc=*/SourceLocation{},
+ /*Id=*/nullptr, FPT->getParamType(i), /*TInfo=*/nullptr,
+ StorageClass::SC_None, /*DefArg=*/nullptr);
+ parm->setScopeInfo(0, i);
+ params.push_back(parm);
+ }
+
+ tempFunc->setParams(params);
+
+ return tempFunc;
+}
+
+std::string
+CIRGenModule::getOpenACCBindMangledName(const IdentifierInfo *bindName,
+ const FunctionDecl *attachedFunction) {
+ FunctionDecl *tempFunc = createOpenACCBindTempFunction(
+ getASTContext(), bindName, attachedFunction);
+
+ std::string ret = getMangledNameImpl(*this, GlobalDecl(tempFunc), tempFunc);
+
+ // This does nothing (it is a do-nothing function), since this is a
+ // slab-allocator, but leave a call in to immediately destroy this in case we
+ // ever come up with a way of getting allocations back.
+ getASTContext().Deallocate(tempFunc);
+ return ret;
+}
+
StringRef CIRGenModule::getMangledName(GlobalDecl gd) {
GlobalDecl canonicalGd = gd.getCanonicalDecl();
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index de263f4868507..2727ca29bf7df 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -507,6 +507,15 @@ class CIRGenModule : public CIRGenTypeCache {
mlir::Value emitMemberPointerConstant(const UnaryOperator *e);
llvm::StringRef getMangledName(clang::GlobalDecl gd);
+ // This function is to support the OpenACC 'bind' clause, which names an
+ // alternate name for the function to be called by. This function mangles
+ // `attachedFunction` as-if its name was actually `bindName` (that is, with
+ // the same signature). It has some additional complications, as the 'bind'
+ // target is always going to be a global function, so member functions need an
+ // explicit instead of implicit 'this' parameter, and thus gets mangled
+ // differently.
+ std::string getOpenACCBindMangledName(const IdentifierInfo *bindName,
+ const FunctionDecl *attachedFunction);
void emitTentativeDefinition(const VarDecl *d);
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-bind.c b/clang/test/CIR/CodeGenOpenACC/routine-bind.c
index 2af024322d67e..72ff952fdd6f4 100644
--- a/clang/test/CIR/CodeGenOpenACC/routine-bind.c
+++ b/clang/test/CIR/CodeGenOpenACC/routine-bind.c
@@ -1,4 +1,5 @@
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+// FIXME: We should run this against Windows mangling as well at one point.
#pragma acc routine seq bind("BIND1")
void Func1(){}
@@ -18,6 +19,28 @@ void Func5(){}
void Func6(){}
#pragma acc routine(Func6) seq device_type(radeon) bind("BIND6_R") device_type(multicore, host) bind("BIND6_M")
+#pragma acc routine seq bind(BIND7)
+void Func7(int i){}
+
+void Func8(float f){}
+#pragma acc routine(Func8) seq bind(BIND8)
+
+#pragma acc routine seq device_type(nvidia) bind(BIND9)
+void Func9(int i, float f, short s){}
+
+struct S{};
+struct U{};
+struct V{};
+
+void Func10(struct S s){}
+#pragma acc routine(Func10) seq device_type(radeon) bind(BIND10)
+
+#pragma acc routine seq device_type(nvidia, host) bind(BIND11_NVH) device_type(multicore) bind(BIND11_MC)
+void Func11(struct U* u, struct V v, int i){}
+
+int Func12(struct U u, struct V v, int i){ return 0; }
+#pragma acc routine(Func12) seq device_type(radeon) bind(BIND12_R) device_type(multicore, host) bind(BIND12_MCH)
+
// 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
//
@@ -33,7 +56,25 @@ void Func6(){}
//
// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
//
+// CHECK: cir.func{{.*}} @[[F7_NAME:.*Func7[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F7_R_NAME]] func(@[[F7_NAME]]) bind(@BIND7) seq
+//
+// CHECK: cir.func{{.*}} @[[F8_NAME:.*Func8[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F9_NAME:.*Func9[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F9_R_NAME]] func(@[[F9_NAME]]) bind(@BIND9 [#acc.device_type<nvidia>]) seq
+//
+// CHECK: cir.func{{.*}} @[[F10_NAME:.*Func10[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F11_NAME:.*Func11[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F11_R_NAME]] func(@[[F11_NAME]]) bind(@BIND11_NVH [#acc.device_type<nvidia>], @BIND11_NVH [#acc.device_type<host>], @BIND11_MC [#acc.device_type<multicore>])
+//
+// CHECK: cir.func{{.*}} @[[F12_NAME:.*Func12[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_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
+// CHECK: acc.routine @[[F8_R_NAME]] func(@[[F8_NAME]]) bind(@BIND8) seq
+// CHECK: acc.routine @[[F10_R_NAME]] func(@[[F10_NAME]]) bind(@BIND10 [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F12_R_NAME]] func(@[[F12_NAME]]) bind(@BIND12_R [#acc.device_type<radeon>], @BIND12_MCH [#acc.device_type<multicore>], @BIND12_MCH [#acc.device_type<host>]) seq
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp b/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
index 2af024322d67e..284196d23376d 100644
--- a/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/routine-bind.cpp
@@ -1,4 +1,5 @@
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+// FIXME: We should run this against Windows mangling as well at one point.
#pragma acc routine seq bind("BIND1")
void Func1(){}
@@ -18,22 +19,126 @@ 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
+#pragma acc routine seq bind(BIND7)
+void Func7(int){}
+
+void Func8(float){}
+#pragma acc routine(Func8) seq bind(BIND8)
+
+#pragma acc routine seq device_type(nvidia) bind(BIND9)
+void Func9(int, float, short){}
+
+struct S{};
+struct U{};
+struct V{};
+
+void Func10(S){}
+#pragma acc routine(Func10) seq device_type(radeon) bind(BIND10)
+
+#pragma acc routine seq device_type(nvidia, host) bind(BIND11_NVH) device_type(multicore) bind(BIND11_MC)
+void Func11(U*, V&, int){}
+
+int Func12(U, V, int){ return 0; }
+#pragma acc routine(Func12) seq device_type(radeon) bind(BIND12_R) device_type(multicore, host) bind(BIND12_MCH)
+
+struct HasFuncs {
+#pragma acc routine seq bind(MEM)
+ int MemFunc(int, double, HasFuncs&, S){ return 0; }
+#pragma acc routine seq bind(MEM)
+ int ConstMemFunc(int, double, HasFuncs&, S) const { return 0; }
+#pragma acc routine seq bind(MEM)
+ int VolatileMemFunc(int, double, HasFuncs&, S) const volatile { return 0; }
+#pragma acc routine seq bind(MEM)
+ int RefMemFunc(int, double, HasFuncs&, S) const && { return 0; }
+#pragma acc routine seq bind(STATICMEM)
+ int StaticMemFunc(int, double, HasFuncs&, U*){ return 0; }
+};
+
+void hasLambdas() {
+ HasFuncs hf;
+ hf.MemFunc(1, 1.0, hf, S{});
+ hf.ConstMemFunc(1, 1.0, hf, S{});
+ static_cast<const volatile HasFuncs>(hf).VolatileMemFunc(1, 1.0, hf, S{});
+ HasFuncs{}.RefMemFunc(1, 1.0, hf, S{});
+ U u;
+ hf.StaticMemFunc(1, 1.0, hf, &u);
+ int i, j, k, l;
+#pragma acc routine seq bind(LAMBDA1)
+ auto Lambda = [](int, float, double){};
+#pragma acc routine seq bind(LAMBDA2)
+ auto Lambda2 = [i, F =&j, k, &l](int, float, double){};
+
+ Lambda(1, 2, 3);
+ Lambda2(1, 2, 3);
+}
+
+// CHECK: cir.func{{.*}} @_Z5Func1v({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F1_R_NAME]] func(@_Z5Func1v) bind("BIND1") seq
//
-// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
+// CHECK: cir.func{{.*}} @_Z5Func2v({{.*}}){{.*}} 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{{.*}} @_Z5Func3v({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F3_R_NAME]] func(@_Z5Func3v) 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{{.*}} @_Z5Func4v({{.*}}){{.*}} 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{{.*}} @_Z5Func5v({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F5_R_NAME]] func(@_Z5Func5v) 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: cir.func{{.*}} @_Z5Func6v({{.*}}){{.*}} 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
+// CHECK: cir.func{{.*}} @_Z5Func7i({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F7_R_NAME]] func(@_Z5Func7i) bind(@_Z5BIND7i) seq
+//
+// CHECK: cir.func{{.*}} @_Z5Func8f({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_Z5Func9ifs({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F9_R_NAME]] func(@_Z5Func9ifs) bind(@_Z5BIND9ifs [#acc.device_type<nvidia>]) seq
+
+// CHECK: cir.func{{.*}} @_Z6Func101S({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_Z6Func11P1UR1Vi({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F11_R_NAME]] func(@_Z6Func11P1UR1Vi) bind(@_Z10BIND11_NVHP1UR1Vi [#acc.device_type<nvidia>], @_Z10BIND11_NVHP1UR1Vi [#acc.device_type<host>], @_Z9BIND11_MCP1UR1Vi [#acc.device_type<multicore>]) seq
+//
+// CHECK: cir.func{{.*}} @_Z6Func121U1Vi({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_ZN8HasFuncs7MemFuncEidRS_1S({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[MEMFUNC_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_ZNK8HasFuncs12ConstMemFuncEidRS_1S({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[CONSTMEMFUNC_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_ZNVK8HasFuncs15VolatileMemFuncEidRS_1S({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[VOLATILEMEMFUNC_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_ZNKO8HasFuncs10RefMemFuncEidRS_1S({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[REFMEMFUNC_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @_ZN8HasFuncs13StaticMemFuncEidRS_P1U({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[STATICFUNC_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} lambda{{.*}} @_ZZ10hasLambdasvENK3$_0clEifd({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[LAMBDA1_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} lambda{{.*}} @_ZZ10hasLambdasvENK3$_1clEifd({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[LAMBDA2_R_NAME:.*]]]>}
+//
+// CHECK: acc.routine @[[MEMFUNC_R_NAME]] func(@_ZN8HasFuncs7MemFuncEidRS_1S) bind(@_Z3MEMP8HasFuncsidRS_1S) seq
+// CHECK: acc.routine @[[CONSTMEMFUNC_R_NAME]] func(@_ZNK8HasFuncs12ConstMemFuncEidRS_1S) bind(@_Z3MEMPK8HasFuncsidRS_1S) seq
+// CHECK: acc.routine @[[VOLATILEMEMFUNC_R_NAME]] func(@_ZNVK8HasFuncs15VolatileMemFuncEidRS_1S) bind(@_Z3MEMPVK8HasFuncsidRS_1S) seq
+// CHECK: acc.routine @[[REFMEMFUNC_R_NAME]] func(@_ZNKO8HasFuncs10RefMemFuncEidRS_1S) bind(@_Z3MEMPK8HasFuncsidRS_1S) seq
+// CHECK: acc.routine @[[STATICFUNC_R_NAME]] func(@_ZN8HasFuncs13StaticMemFuncEidRS_P1U) bind(@_Z9STATICMEMP8HasFuncsidRS_P1U) seq
+//
+// These two LOOK weird because the first argument to each of these is the
+// implicit 'this', so they look like they have the lambda mangling (and
+// demanglers don't handle lambdas well).
+// CHECK: acc.routine @[[LAMBDA1_R_NAME]] func(@_ZZ10hasLambdasvENK3$_0clEifd) bind(@_Z7LAMBDA1PKZ10hasLambdasvE3$_0ifd) seq
+// Manual demangle:
+// Func name: _Z7LAMBDA1 -> LAMBDA1
+// Args: P -> Pointer
+// K -> Const
+// Z10hasLambdasv-> hasLambdas(void)::
+// E3$_0 -> anonymous type #0
+// ifd -> taking args int, float, double.
+// // CHECK: acc.routine @[[LAMBDA2_R_NAME]] func(@_ZZ10hasLambdasvENK3$_1clEifd) bind(@_Z7LAMBDA2PKZ10hasLambdasvE3$_1ifd) seq
+
+// CHECK: acc.routine @[[F2_R_NAME]] func(@_Z5Func2v) bind("BIND2") seq
+// CHECK: acc.routine @[[F4_R_NAME]] func(@_Z5Func4v) bind("BIND4" [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F6_R_NAME]] func(@_Z5Func6v) bind("BIND6_R" [#acc.device_type<radeon>], "BIND6_M" [#acc.device_type<multicore>], "BIND6_M" [#acc.device_type<host>]) seq
+// CHECK: acc.routine @[[F8_R_NAME]] func(@_Z5Func8f) bind(@_Z5BIND8f) seq
+// CHECK: acc.routine @[[F10_R_NAME]] func(@_Z6Func101S) bind(@_Z6BIND101S [#acc.device_type<radeon>]) seq
+// CHECK: acc.routine @[[F12_R_NAME]] func(@_Z6Func121U1Vi) bind(@_Z8BIND12_R1U1Vi [#acc.device_type<radeon>], @_Z10BIND12_MCH1U1Vi [#acc.device_type<multicore>], @_Z10BIND12_MCH1U1Vi [#acc.device_type<host>]) seq
|
andykaylor
approved these changes
Dec 11, 2025
Contributor
andykaylor
left a comment
andykaylor
left a comment
There was a problem hiding this comment.
Just a couple of nits from me. I'll leave it to @razvanlupusoru to comment on your interpretation of the standard.
| protoFunc->getType(), /*TInfo=*/nullptr, StorageClass::SC_None); | ||
|
|
||
| QualType funcTy = protoFunc->getType(); | ||
| auto *FPT = cast<FunctionProtoType>(protoFunc->getType()); |
Contributor
There was a problem hiding this comment.
Suggested change
| auto *FPT = cast<FunctionProtoType>(protoFunc->getType()); | |
| auto *fpt = cast<FunctionProtoType>(protoFunc->getType()); |
Collaborator
Author
There was a problem hiding this comment.
Urgh, almost made it...
| /*NLoc=*/SourceLocation{}, bindName, funcTy, | ||
| /*TInfo=*/nullptr, StorageClass::SC_None); | ||
|
|
||
| SmallVector<ParmVarDecl *, 16> params; |
Contributor
There was a problem hiding this comment.
Why 16? Can you drop that and reserve instead?
Collaborator
Author
There was a problem hiding this comment.
Just a copy/paste from the function-emit code, I didn't put much thought into it. A reserve seems good enough, I'll do that.
Contributor
This seems sensible to me and glad you found a reasonable way to do this. |
razvanlupusoru
approved these changes
Dec 12, 2025
anonymouspc
pushed a commit
to anonymouspc/llvm
that referenced
this pull request
Dec 15, 2025
The bind clause specifies the name of the function to call on the device, and takes either a string or identifier(per the standard): "If the name is specified as an identifier, it is callled as if the name were specified in the language being compiled. If the name is specified as a string, the string is used for the procedure name unmodified". The latter (as a string) is already implemented, this patch implements the former. Unfortunately, no existing implementation of this in C++ seems to exist. Other languages, the 'name' of a function is sufficient to identify it (in this case 'bind' can refer to undeclared functions), so it is possible to figure out what the name should be. In C++ with overloading (without a discriminator, ala-fortran), a name only names an infinite overload set. SO, in order to implement this, I've decided that the 'called as' (bound) function must have the same signature as the one marked by the 'routine'. This is trivially sensible in non-member functions, however requires a bit more thought for member(and thus lambda-call-operators) functions. In this case, we 'promote' the type of the function to a 'free' function by turning the implicit 'this' to an explicit 'this'. I believe this is the most sensible and reasonable way to implement this, and really the only way to make something usable.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The bind clause specifies the name of the function to call on the device, and takes either a string or identifier(per the standard):
"If the name is specified as an identifier, it is callled as if the name were specified in the language being compiled. If the name is specified as a string, the string is used for the procedure name unmodified".
The latter (as a string) is already implemented, this patch implements the former. Unfortunately, no existing implementation of this in C++ seems to exist. Other languages, the 'name' of a function is sufficient to identify it (in this case 'bind' can refer to undeclared functions), so it is possible to figure out what the name should be. In C++ with overloading (without a discriminator, ala-fortran), a name only names an infinite overload set.
SO, in order to implement this, I've decided that the 'called as' (bound) function must have the same signature as the one marked by the 'routine'. This is trivially sensible in non-member functions, however requires a bit more thought for member(and thus lambda-call-operators) functions. In this case, we 'promote' the type of the function to a 'free' function by turning the implicit 'this' to an explicit 'this'.
I believe this is the most sensible and reasonable way to implement this, and really the only way to make something usable.