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

[scott.linder]

Summary

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.

[scott.linder]

I'm not certain I fully understand NVPTX's relationship with its debugger, but from https://reviews.llvm.org/D57162 I gather the "default" address space in the debugger is global, and so the frontend omits it rather than explicitly mentioning it. I think it would be simpler to carry this information throughout the compiler, and only strip it late in the backend as a quirk controllable via some "optimize for NVPTX debugger", but in the patch as it currently is I instead just update NVPTXDWARFAddrSpaceMap.

Edit: Concerning auto variables, when coming back to the patch to post it I had missed the next patch in the series which addresses it by directly referring to the corresponding alloca rather than the addrspacecast to the default address space. I'll post that patch shortly to address the "FIXME" in this one.

[aprantl]

That looks much nicer.

clang/test/CodeGenHIP/debug-info-address-class.hip

9

They are more convenient, but having very many CHECK_DAGs is also really slow — would it be feasible to reorder them and use CHECKs? Perhaps by running clang/FileCheck twice with different sets of CHECK lines?

[scott.linder]

Replace uses of CHECK-DAG, use more meaningful names in test

[scott.linder]

clang/test/CodeGenHIP/debug-info-address-class.hip

9

The test is pretty short, so I just re-ordered the checks to match how they appear in the output (and used more descriptive names to make it easier to follow).

This does mean the test relies on the order these things are traversed. Some bits are maybe a bit surprising, like how the metadata for the __shared__ auto variable comes before the argument, but I don't imagine it is liable to change often/accidentally.

[scott.linder]

clang/lib/Basic/Targets/NVPTX.h

47

Does anyone have any thoughts on this change specifically? Is someone more familiar with NVPTX willing to weigh in on whether it makes more sense to carry the address space throughout the compiler explicitly and "drop" it late in the DWARF emission, or to do what I did in the current patch (drop it early).

I would lean towards updating the patch to do the latter, but I wanted to get feedback before plunging off to do it.

[arsenm]

ping?

[yaxunl]

LGTM. Thanks.

clang/lib/Basic/Targets/NVPTX.h

 //===--- NVPTX.h - Declare NVPTX target feature support ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file declares NVPTX TargetInfo objects.
static const unsigned NVPTXAddrSpaceMap[] = {
     0, // ptr32_uptr
     0  // ptr64
 };
 
 /// The DWARF address class. Taken from
 /// 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
 };
 
 class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
   static const char *const GCCRegNames[];
   static const Builtin::Info BuiltinInfo[];
   CudaArch GPU;
   uint32_t PTXVersion;

clang/lib/CodeGen/CGDebugInfo.cpp

if (T->isUnionType() && DeclName.empty()) {
     assert(RD->isAnonymousStructOrUnion() &&
            "unnamed non-anonymous struct or union?");
     GVE = CollectAnonRecordDecls(RD, Unit, LineNo, LinkageName, Var, DContext);
   } else {
     auto Align = getDeclAlignIfRequired(D, CGM.getContext());
 
     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(
         DContext, DeclName, LinkageName, Unit, LineNo, getOrCreateType(T, Unit),
         Var->hasLocalLinkage(), true,
         Expr.empty() ? nullptr : DBuilder.createExpression(Expr),
         getOrCreateStaticDataMemberDeclarationOrNull(D), TemplateParameters,
         Align);

clang/test/CodeGenHIP/debug-info-address-class.hip

+// 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))
+
+__device__ int FileVarDevice;
+__device__ __shared__ int FileVarDeviceShared;
+__device__ __constant__ int FileVarDeviceConstant;
+
+__device__ void kernel1(
+    // FIXME This should be in the private address space.
+    // CHECK: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[ARG:[0-9]+]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+    int Arg) {
+  __shared__ int FuncVarShared;
+
+  // FIXME This should be in the private address space.
+  // CHECK: call void @llvm.dbg.declare(metadata i32* {{.*}}, metadata ![[FUNC_VAR:[0-9]+]], metadata !DIExpression()), !dbg !{{[0-9]+}}
+  int FuncVar;
+}
+
+// CHECK: !DIGlobalVariableExpression(var: ![[FILE_VAR_DEVICE:[0-9]+]], expr: !DIExpression())
+// CHECK: ![[FILE_VAR_DEVICE]] = distinct !DIGlobalVariable(name: "FileVarDevice", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+
+// CHECK: !DIGlobalVariableExpression(var: ![[FILE_VAR_DEVICE_SHARED:[0-9]+]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
+// CHECK: ![[FILE_VAR_DEVICE_SHARED]] = distinct !DIGlobalVariable(name: "FileVarDeviceShared", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+
+// CHECK: !DIGlobalVariableExpression(var: ![[FILE_VAR_DEVICE_CONSTANT:[0-9]+]], expr: !DIExpression())
+// CHECK: ![[FILE_VAR_DEVICE_CONSTANT]] = distinct !DIGlobalVariable(name: "FileVarDeviceConstant", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true)
+
+// CHECK: !DIGlobalVariableExpression(var: ![[FUNC_VAR_SHARED:[0-9]+]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
+// CHECK: ![[FUNC_VAR_SHARED]] = distinct !DIGlobalVariable(name: "FuncVarShared", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true)
+
+// CHECK: ![[ARG]] = !DILocalVariable(name: "Arg", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
+
+// CHECK: ![[FUNC_VAR]] = !DILocalVariable(name: "FuncVar", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})