[CUDA][HIP][DebugInfo] Skip reference device function

[hliao]

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.

[hliao]

make the test more robust to potential metadata identifier change

[aprantl]

clang/lib/CodeGen/CGDebugInfo.cpp

1755

This wasn't there before... why is this necessary?

[hliao]

clang/lib/CodeGen/CGDebugInfo.cpp

1755

That's from the original line 1752.

[aprantl]

clang/lib/CodeGen/CGDebugInfo.cpp

1755

Oh :-)

clang/lib/CodeGen/CGDebugInfo.cpp

case TemplateArgument::Integral: {
           TheCU, Name, TTy,
           llvm::ConstantInt::get(CGM.getLLVMContext(), TA.getAsIntegral())));
     } break;
     case TemplateArgument::Declaration: {
       const ValueDecl *D = TA.getAsDecl();
       QualType T = TA.getParamTypeForDecl().getDesugaredType(CGM.getContext());
       llvm::DIType *TTy = getOrCreateType(T, Unit);
       llvm::Constant *V = nullptr;
+      // 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.
+        // 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())) {
+        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();
       llvm::DIType *TTy = getOrCreateType(T, Unit);
       llvm::Constant *V = nullptr;
       // Special case member data pointer null values since they're actually -1
       // instead of zero.
       if (const auto *MPT = dyn_cast<MemberPointerType>(T.getTypePtr()))

clang/test/CodeGenCUDA/debug-info-template.cu

+// 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)