Automerge: [OpenACC][CIR] 'bind' lowering with identifier (#171749)

erichkeane · github-actions[bot] · commit 5e2bf0161f5f · 2025-12-12T14:15:43.000Z
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.
diff --git 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
@@ -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
@@ -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
@@ -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
@@ -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
 
