[PATCH] D88976: [clang] Use correct address space for global variable debug info

[Scott Linder via Phabricator via cfe-commits]

scott.linder created this revision.
Herald added subscribers: cfe-commits, jholewinski.
Herald added a project: clang.
scott.linder requested review of this revision.

The target needs to be queried here, but previously we seemed to only
duplicate CUDA's (and so HIP's) behavior, and only partially. Use the
same function as codegen to find the correct address space.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D88976

Files:
  clang/lib/Basic/Targets/NVPTX.h
  clang/lib/CodeGen/CGDebugInfo.cpp
  clang/test/CodeGenHIP/debug-info-address-class.hip


Index: clang/test/CodeGenHIP/debug-info-address-class.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGenHIP/debug-info-address-class.hip
@@ -0,0 +1,31 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -dwarf-version=4 -o - %s | FileCheck %s
+
+#define __device__ __attribute__((device))
+#define __shared__ __attribute__((shared))
+#define __constant__ __attribute__((constant))
+
+// CHECK-DAG: ![[FILEVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR0]], expr: !DIExpression())
+__device__ int FileVar0;
+// CHECK-DAG: ![[FILEVAR1:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR1]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
+__device__ __shared__ int FileVar1;
+// CHECK-DAG: ![[FILEVAR2:[0-9]+]] = distinct !DIGlobalVariable(name: "FileVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+// CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR2]], expr: !DIExpression())
+__device__ __constant__ int FileVar2;
+
+__device__ void kernel1(
+    // FIXME This should be in the private address space.
+    // CHECK-DAG: ![[ARG:[0-9]+]] = !DILocalVariable(name: "Arg", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
+    // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[ARG]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+    int Arg) {
+    // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true)
+    // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR0]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
+  __shared__ int FuncVar0;
+
+  // FIXME This should be in the private address space.
+  // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
+  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+  int FuncVar1;
+}
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -4664,15 +4664,7 @@
 
     SmallVector<int64_t, 4> Expr;
     unsigned AddressSpace =
-        CGM.getContext().getTargetAddressSpace(D->getType());
-    if (CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) {
-      if (D->hasAttr<CUDASharedAttr>())
-        AddressSpace =
-            CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared);
-      else if (D->hasAttr<CUDAConstantAttr>())
-        AddressSpace =
-            CGM.getContext().getTargetAddressSpace(LangAS::cuda_constant);
-    }
+        CGM.getContext().getTargetAddressSpace(CGM.GetGlobalVarAddressSpace(D));
     AppendAddressSpaceXDeref(AddressSpace, Expr);
 
     GVE = DBuilder.createGlobalVariableExpression(
Index: clang/lib/Basic/Targets/NVPTX.h
===================================================================
--- clang/lib/Basic/Targets/NVPTX.h
+++ clang/lib/Basic/Targets/NVPTX.h
@@ -44,7 +44,7 @@
 /// https://docs.nvidia.com/cuda/archive/10.0/ptx-writers-guide-to-interoperability/index.html#cuda-specific-dwarf
 static const int NVPTXDWARFAddrSpaceMap[] = {
     -1, // Default, opencl_private or opencl_generic - not defined
-    5,  // opencl_global
+    -1, // opencl_global
     -1,
     8,  // opencl_local or cuda_shared
     4,  // opencl_constant or cuda_constant


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D88976.296705.patch
Type: text/x-patch
Size: 4101 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201007/40b65ed7/attachment.bin>