Author: hliao Date: Wed Mar 6 13:16:27 2019 New Revision: 355551 URL: http://llvm.org/viewvc/llvm-project?rev=355551&view=rev Log: [CUDA][HIP][DebugInfo] Skip reference device function
Summary: - A device functions could be used as a non-type template parameter in a global/host function template. However, we should not try to retrieve that device function and reference it in the host-side debug info as it's only valid at device side. Subscribers: aprantl, jdoerfert, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D58992 Added: cfe/trunk/test/CodeGenCUDA/debug-info-template.cu Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDebugInfo.cpp?rev=355551&r1=355550&r2=355551&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp (original) +++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Wed Mar 6 13:16:27 2019 @@ -1725,31 +1725,37 @@ CGDebugInfo::CollectTemplateParams(const QualType T = TA.getParamTypeForDecl().getDesugaredType(CGM.getContext()); llvm::DIType *TTy = getOrCreateType(T, Unit); llvm::Constant *V = nullptr; - const CXXMethodDecl *MD; - // Variable pointer template parameters have a value that is the address - // of the variable. - if (const auto *VD = dyn_cast<VarDecl>(D)) - V = CGM.GetAddrOfGlobalVar(VD); - // Member function pointers have special support for building them, though - // this is currently unsupported in LLVM CodeGen. - else if ((MD = dyn_cast<CXXMethodDecl>(D)) && MD->isInstance()) - V = CGM.getCXXABI().EmitMemberFunctionPointer(MD); - else if (const auto *FD = dyn_cast<FunctionDecl>(D)) - V = CGM.GetAddrOfFunction(FD); - // Member data pointers have special handling too to compute the fixed - // offset within the object. - else if (const auto *MPT = dyn_cast<MemberPointerType>(T.getTypePtr())) { - // These five lines (& possibly the above member function pointer - // handling) might be able to be refactored to use similar code in - // CodeGenModule::getMemberPointerConstant - uint64_t fieldOffset = CGM.getContext().getFieldOffset(D); - CharUnits chars = - CGM.getContext().toCharUnitsFromBits((int64_t)fieldOffset); - V = CGM.getCXXABI().EmitMemberDataPointer(MPT, chars); + // Skip retrieve the value if that template parameter has cuda device + // attribute, i.e. that value is not available at the host side. + if (!CGM.getLangOpts().CUDA || CGM.getLangOpts().CUDAIsDevice || + !D->hasAttr<CUDADeviceAttr>()) { + const CXXMethodDecl *MD; + // Variable pointer template parameters have a value that is the address + // of the variable. + if (const auto *VD = dyn_cast<VarDecl>(D)) + V = CGM.GetAddrOfGlobalVar(VD); + // Member function pointers have special support for building them, + // though this is currently unsupported in LLVM CodeGen. + else if ((MD = dyn_cast<CXXMethodDecl>(D)) && MD->isInstance()) + V = CGM.getCXXABI().EmitMemberFunctionPointer(MD); + else if (const auto *FD = dyn_cast<FunctionDecl>(D)) + V = CGM.GetAddrOfFunction(FD); + // Member data pointers have special handling too to compute the fixed + // offset within the object. + else if (const auto *MPT = + dyn_cast<MemberPointerType>(T.getTypePtr())) { + // These five lines (& possibly the above member function pointer + // handling) might be able to be refactored to use similar code in + // CodeGenModule::getMemberPointerConstant + uint64_t fieldOffset = CGM.getContext().getFieldOffset(D); + CharUnits chars = + CGM.getContext().toCharUnitsFromBits((int64_t)fieldOffset); + V = CGM.getCXXABI().EmitMemberDataPointer(MPT, chars); + } + V = V->stripPointerCasts(); } TemplateParams.push_back(DBuilder.createTemplateValueParameter( - TheCU, Name, TTy, - cast_or_null<llvm::Constant>(V->stripPointerCasts()))); + TheCU, Name, TTy, cast_or_null<llvm::Constant>(V))); } break; case TemplateArgument::NullPtr: { QualType T = TA.getNullPtrType(); Added: cfe/trunk/test/CodeGenCUDA/debug-info-template.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/debug-info-template.cu?rev=355551&view=auto ============================================================================== --- cfe/trunk/test/CodeGenCUDA/debug-info-template.cu (added) +++ cfe/trunk/test/CodeGenCUDA/debug-info-template.cu Wed Mar 6 13:16:27 2019 @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - -debug-info-kind=limited -dwarf-version=2 -debugger-tuning=gdb | FileCheck %s + +#include "Inputs/cuda.h" + +__device__ void f(); +template<void(*F)()> __global__ void t() { F(); } +__host__ void g() { t<f><<<1,1>>>(); } + +// Ensure the value of device-function (as value template parameter) is null. +// CHECK: !DITemplateValueParameter(name: "F", type: !{{[0-9]+}}, value: null) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
