#include "clang/AST/DeclCXX.h"
#include "clang/AST/DeclObjC.h"
#include "clang/AST/DeclOpenMP.h"
+#include "clang/AST/DeclTemplate.h"
#include "clang/Basic/ABI.h"
#include "clang/Basic/LLVM.h"
#include "llvm/ADT/DenseMapInfo.h"
}
KernelReferenceKind getKernelReferenceKind() const {
- assert(isa<FunctionDecl>(getDecl()) &&
- cast<FunctionDecl>(getDecl())->hasAttr<CUDAGlobalAttr>() &&
+ assert(((isa<FunctionDecl>(getDecl()) &&
+ cast<FunctionDecl>(getDecl())->hasAttr<CUDAGlobalAttr>()) ||
+ (isa<FunctionTemplateDecl>(getDecl()) &&
+ cast<FunctionTemplateDecl>(getDecl())
+ ->getTemplatedDecl()
+ ->hasAttr<CUDAGlobalAttr>())) &&
"Decl is not a GPU kernel!");
return static_cast<KernelReferenceKind>(Value.getInt());
}
switch (Name.getNameKind()) {
case DeclarationName::Identifier: {
if (const IdentifierInfo *II = Name.getAsIdentifierInfo()) {
- mangleSourceName(II->getName());
+ bool IsDeviceStub =
+ ND &&
+ ((isa<FunctionDecl>(ND) && ND->hasAttr<CUDAGlobalAttr>()) ||
+ (isa<FunctionTemplateDecl>(ND) &&
+ cast<FunctionTemplateDecl>(ND)
+ ->getTemplatedDecl()
+ ->hasAttr<CUDAGlobalAttr>())) &&
+ GD.getKernelReferenceKind() == KernelReferenceKind::Stub;
+ if (IsDeviceStub)
+ mangleSourceName(
+ (llvm::Twine("__device_stub__") + II->getName()).str());
+ else
+ mangleSourceName(II->getName());
break;
}
// GNU: @[[HDKERN:_Z11kernel_declv]] = external constant void ()*, align 8
// MSVC: @[[HCKERN:ckernel]] = dso_local constant void ()* @[[CSTUB:__device_stub__ckernel]], align 8
-// MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?nskernel@ns@@YAXXZ"]], align 8
-// MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$kernelfunc@H@@YAXXZ.*"]], comdat, align 8
+// MSVC: @[[HNSKERN:"\?nskernel@ns@@YAXXZ.*"]] = dso_local constant void ()* @[[NSSTUB:"\?__device_stub__nskernel@ns@@YAXXZ"]], align 8
+// MSVC: @[[HTKERN:"\?\?\$kernelfunc@H@@YAXXZ.*"]] = linkonce_odr dso_local constant void ()* @[[TSTUB:"\?\?\$__device_stub__kernelfunc@H@@YAXXZ.*"]], comdat, align 8
// MSVC: @[[HDKERN:"\?kernel_decl@@YAXXZ.*"]] = external dso_local constant void ()*, align 8
extern "C" __global__ void ckernel() {}
// CHECK: call void @[[NSSTUB]]()
// CHECK: call void @[[TSTUB]]()
// GNU: call void @[[DSTUB:_Z26__device_stub__kernel_declv]]()
-// MSVC: call void @[[DSTUB:"\?kernel_decl@@YAXXZ"]]()
+// MSVC: call void @[[DSTUB:"\?__device_stub__kernel_decl@@YAXXZ"]]()
extern "C" void fun1(void) {
ckernel<<<1, 1>>>();