[PATCH] D58992: [CUDA][HIP][DebugInfo] Skip reference device function

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
Herald added subscribers: cfe-commits, jdoerfert, aprantl.
Herald added a project: clang.

- 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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D58992

Files:
  clang/lib/CodeGen/CGDebugInfo.cpp
  clang/test/CodeGenCUDA/debug-info-template.cu


Index: clang/test/CodeGenCUDA/debug-info-template.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/debug-info-template.cu
@@ -0,0 +1,11 @@
+// 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) in the is
+// null.
+// CHECK: !DITemplateValueParameter(name: "F", type: !34, value: null)
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -1725,31 +1725,37 @@
       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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D58992.189383.patch
Type: text/x-patch
Size: 4217 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190305/e717d007/attachment.bin>