Author: Erich Keane Date: 2025-12-12T06:04:47-08:00 New Revision: 1dbff713127ad97c118a9ce64b12d014d3252b01
URL: https://github.com/llvm/llvm-project/commit/1dbff713127ad97c118a9ce64b12d014d3252b01 DIFF: https://github.com/llvm/llvm-project/commit/1dbff713127ad97c118a9ce64b12d014d3252b01.diff LOG: [OpenACC][CIR] 'bind' lowering with identifier (#171749) 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. Added: Modified: clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp clang/lib/CIR/CodeGen/CIRGenModule.cpp clang/lib/CIR/CodeGen/CIRGenModule.h clang/test/CIR/CodeGenOpenACC/routine-bind.c clang/test/CIR/CodeGenOpenACC/routine-bind.cpp Removed: ################################################################################ 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 edaf1c46eb097..9564c9efbc9a0 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -1721,6 +1721,71 @@ 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 *> params; + params.reserve(fpt->getNumParams()); + + // 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 + // diff erently. + 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
