merge amd-debug into amd-staging (llvm#3753)

Author: ronlieb

clang/lib/Basic/Targets/AMDGPU.h

@@ -400,15 +400,12 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
   /// in the DWARF.
   std::optional<unsigned>
   getDWARFAddressSpace(unsigned AddressSpace) const override {
-    const unsigned DWARF_Private = 1;
-    const unsigned DWARF_Local = 2;
-    if (AddressSpace == llvm::AMDGPUAS::PRIVATE_ADDRESS) {
-      return DWARF_Private;
-    } else if (AddressSpace == llvm::AMDGPUAS::LOCAL_ADDRESS) {
-      return DWARF_Local;
-    } else {
+    int DWARFAS = llvm::AMDGPU::mapToDWARFAddrSpace(AddressSpace);
+    // If there is no corresponding address space identifier, or it would be
+    // the default, then don't emit the attribute.
+    if (DWARFAS == -1 || DWARFAS == llvm::AMDGPU::DWARFAS::DEFAULT)
       return std::nullopt;
-    }
+    return DWARFAS;
   }
 
   CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -5436,7 +5436,7 @@ llvm::DILocalVariable *CGDebugInfo::EmitDeclareForHeterogeneousDwarf(
   } else if (const auto *RT = dyn_cast<RecordType>(VD->getType())) {
     // If VD is an anonymous union then Storage represents value for
     // all union fields.
-    const RecordDecl *RD = RT->getOriginalDecl();
+    const RecordDecl *RD = RT->getOriginalDecl()->getDefinitionOrSelf();
     if (RD->isUnion() && RD->isAnonymousStructOrUnion()) {
       llvm::DIExprBuilder UnionExprBuilder{ExprBuilder};
       llvm::DIExpression *UnionDIExpression = UnionExprBuilder.intoExpression();
@@ -5999,9 +5999,9 @@ llvm::DIGlobalVariableExpression *CGDebugInfo::CollectAnonRecordDecls(
     // Ignore unnamed fields, but recurse into anonymous records.
     if (FieldName.empty()) {
       if (const auto *RT = dyn_cast<RecordType>(Field->getType()))
-        GVE =
-            CollectAnonRecordDecls(RT->getOriginalDecl()->getDefinitionOrSelf(),
-                                   Unit, LineNo, LinkageName, MS, Var, DContext);
+        GVE = CollectAnonRecordDecls(
+            RT->getOriginalDecl()->getDefinitionOrSelf(), Unit, LineNo,
+            LinkageName, MS, Var, DContext);
       continue;
     }
     // Use VarDecl's Tag, Scope and Line number.
@@ -6030,7 +6030,8 @@ CGDebugInfo::CollectAnonRecordDeclsForHeterogeneousDwarf(
     if (FieldName.empty()) {
       if (const auto *RT = dyn_cast<RecordType>(Field->getType()))
         GVE = CollectAnonRecordDeclsForHeterogeneousDwarf(
-            RT->getOriginalDecl(), Unit, LineNo, LinkageName, MS, Var, DContext);
+            RT->getOriginalDecl()->getDefinitionOrSelf(), Unit, LineNo,
+            LinkageName, MS, Var, DContext);
       continue;
     }
     // Use VarDecl's Tag, Scope and Line number.
@@ -6393,7 +6394,8 @@ void CGDebugInfo::EmitGlobalVariableForHeterogeneousDwarf(
   // to find the name of any field in the union.
   llvm::dwarf::MemorySpace MS = getDWARFMemorySpace(D);
   if (T->isUnionType() && DeclName.empty()) {
-    const RecordDecl *RD = T->castAs<RecordType>()->getOriginalDecl();
+    const RecordDecl *RD =
+        T->castAs<RecordType>()->getOriginalDecl()->getDefinitionOrSelf();
     assert(RD->isAnonymousStructOrUnion() &&
            "unnamed non-anonymous struct or union?");
     // FIXME(KZHURAVL): No tests for this path.
@@ -6514,7 +6516,6 @@ void CGDebugInfo::EmitGlobalVariableForHeterogeneousDwarf(
 
   if (const auto *ECD = dyn_cast<EnumConstantDecl>(VD)) {
     const auto *ED = cast<EnumDecl>(ECD->getDeclContext());
-    assert(isa<EnumType>(ED->getASTContext().getCanonicalTagType(ED)) && "Enum without EnumType?");
 
     if (CGM.getCodeGenOpts().EmitCodeView) {
       // If CodeView, emit enums as global variables, unless they are defined
@@ -6527,10 +6528,9 @@ void CGDebugInfo::EmitGlobalVariableForHeterogeneousDwarf(
       // If not CodeView, emit DW_TAG_enumeration_type if necessary. For
       // example: for "enum { ZERO };", a DW_TAG_enumeration_type is created the
       // first time `ZERO` is referenced in a function.
-      llvm::DIType *EDTy =
-          getOrCreateType(QualType(ED->getASTContext().getCanonicalTagType(ED)), Unit);
-      assert (EDTy->getTag() == llvm::dwarf::DW_TAG_enumeration_type);
-      (void)EDTy;
+      CanQualType T = CGM.getContext().getCanonicalTagType(ED);
+      [[maybe_unused]] llvm::DIType *EDTy = getOrCreateType(T, Unit);
+      assert(EDTy->getTag() == llvm::dwarf::DW_TAG_enumeration_type);
       return;
     }
   }

clang/test/CodeGenHIP/debug-info-amdgcn-abi-heterogeneous-dwarf.hip

@@ -1096,13 +1096,13 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[META75]] = !DISubprogram(name: "StructTrivialCopyTrivialMove", scope: [[META72]], file: [[META4]], line: 41, type: [[META76:![0-9]+]], scopeLine: 41, flags: DIFlagPrototyped, spFlags: 0)
 // CHECK: [[META76]] = !DISubroutineType(types: [[META77:![0-9]+]])
 // CHECK: [[META77]] = !{null, [[META78:![0-9]+]], [[META79:![0-9]+]]}
-// CHECK: [[META78]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META72]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer)
-// CHECK: [[META79]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META80:![0-9]+]], size: 64)
+// CHECK: [[META78]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META72]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer, addressSpace: 1)
+// CHECK: [[META79]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META80:![0-9]+]], size: 64, addressSpace: 1)
 // CHECK: [[META80]] = !DIDerivedType(tag: DW_TAG_const_type, baseType: [[META72]])
 // CHECK: [[META81]] = !DISubprogram(name: "StructTrivialCopyTrivialMove", scope: [[META72]], file: [[META4]], line: 42, type: [[META82:![0-9]+]], scopeLine: 42, flags: DIFlagPrototyped, spFlags: 0)
 // CHECK: [[META82]] = !DISubroutineType(types: [[META83:![0-9]+]])
 // CHECK: [[META83]] = !{null, [[META78]], [[META84:![0-9]+]]}
-// CHECK: [[META84]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META72]], size: 64)
+// CHECK: [[META84]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META72]], size: 64, addressSpace: 1)
 // CHECK: [[META85]] = !{[[META86]]}
 // CHECK: [[META86]] = !DILocalVariable(arg: 1, scope: [[DBG69]], file: [[META4]], line: 160, type: [[META72]])
 // CHECK: [[META87]] = !DILocation(line: 160, column: 84, scope: [[DBG69]])
@@ -1121,13 +1121,13 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[META100]] = !DISubprogram(name: "StructNoCopyTrivialMove", scope: [[META97]], file: [[META4]], line: 46, type: [[META101:![0-9]+]], scopeLine: 46, flags: DIFlagPrototyped, spFlags: DISPFlagDeleted)
 // CHECK: [[META101]] = !DISubroutineType(types: [[META102:![0-9]+]])
 // CHECK: [[META102]] = !{null, [[META103:![0-9]+]], [[META104:![0-9]+]]}
-// CHECK: [[META103]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META97]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer)
-// CHECK: [[META104]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META105:![0-9]+]], size: 64)
+// CHECK: [[META103]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META97]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer, addressSpace: 1)
+// CHECK: [[META104]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META105:![0-9]+]], size: 64, addressSpace: 1)
 // CHECK: [[META105]] = !DIDerivedType(tag: DW_TAG_const_type, baseType: [[META97]])
 // CHECK: [[META106]] = !DISubprogram(name: "StructNoCopyTrivialMove", scope: [[META97]], file: [[META4]], line: 47, type: [[META107:![0-9]+]], scopeLine: 47, flags: DIFlagPrototyped, spFlags: 0)
 // CHECK: [[META107]] = !DISubroutineType(types: [[META108:![0-9]+]])
 // CHECK: [[META108]] = !{null, [[META103]], [[META109:![0-9]+]]}
-// CHECK: [[META109]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META97]], size: 64)
+// CHECK: [[META109]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META97]], size: 64, addressSpace: 1)
 // CHECK: [[META110]] = !{[[META111]]}
 // CHECK: [[META111]] = !DILocalVariable(arg: 1, scope: [[DBG94]], file: [[META4]], line: 182, type: [[META97]])
 // CHECK: [[META112]] = !DILocation(line: 182, column: 74, scope: [[DBG94]])
@@ -1146,13 +1146,13 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[META125]] = !DISubprogram(name: "StructTrivialCopyNoMove", scope: [[META122]], file: [[META4]], line: 51, type: [[META126:![0-9]+]], scopeLine: 51, flags: DIFlagPrototyped, spFlags: 0)
 // CHECK: [[META126]] = !DISubroutineType(types: [[META127:![0-9]+]])
 // CHECK: [[META127]] = !{null, [[META128:![0-9]+]], [[META129:![0-9]+]]}
-// CHECK: [[META128]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META122]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer)
-// CHECK: [[META129]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META130:![0-9]+]], size: 64)
+// CHECK: [[META128]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META122]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer, addressSpace: 1)
+// CHECK: [[META129]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META130:![0-9]+]], size: 64, addressSpace: 1)
 // CHECK: [[META130]] = !DIDerivedType(tag: DW_TAG_const_type, baseType: [[META122]])
 // CHECK: [[META131]] = !DISubprogram(name: "StructTrivialCopyNoMove", scope: [[META122]], file: [[META4]], line: 52, type: [[META132:![0-9]+]], scopeLine: 52, flags: DIFlagPrototyped, spFlags: DISPFlagDeleted)
 // CHECK: [[META132]] = !DISubroutineType(types: [[META133:![0-9]+]])
 // CHECK: [[META133]] = !{null, [[META128]], [[META134:![0-9]+]]}
-// CHECK: [[META134]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META122]], size: 64)
+// CHECK: [[META134]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META122]], size: 64, addressSpace: 1)
 // CHECK: [[META135]] = !{[[META136]]}
 // CHECK: [[META136]] = !DILocalVariable(arg: 1, scope: [[DBG119]], file: [[META4]], line: 204, type: [[META122]])
 // CHECK: [[META137]] = !DILocation(line: 204, column: 74, scope: [[DBG119]])
@@ -1171,13 +1171,13 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[META150]] = !DISubprogram(name: "StructNoCopyNoMove", scope: [[META147]], file: [[META4]], line: 56, type: [[META151:![0-9]+]], scopeLine: 56, flags: DIFlagPrototyped, spFlags: DISPFlagDeleted)
 // CHECK: [[META151]] = !DISubroutineType(types: [[META152:![0-9]+]])
 // CHECK: [[META152]] = !{null, [[META153:![0-9]+]], [[META154:![0-9]+]]}
-// CHECK: [[META153]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META147]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer)
-// CHECK: [[META154]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META155:![0-9]+]], size: 64)
+// CHECK: [[META153]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META147]], size: 64, flags: DIFlagArtificial | DIFlagObjectPointer, addressSpace: 1)
+// CHECK: [[META154]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META155:![0-9]+]], size: 64, addressSpace: 1)
 // CHECK: [[META155]] = !DIDerivedType(tag: DW_TAG_const_type, baseType: [[META147]])
 // CHECK: [[META156]] = !DISubprogram(name: "StructNoCopyNoMove", scope: [[META147]], file: [[META4]], line: 57, type: [[META157:![0-9]+]], scopeLine: 57, flags: DIFlagPrototyped, spFlags: DISPFlagDeleted)
 // CHECK: [[META157]] = !DISubroutineType(types: [[META158:![0-9]+]])
 // CHECK: [[META158]] = !{null, [[META153]], [[META159:![0-9]+]]}
-// CHECK: [[META159]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META147]], size: 64)
+// CHECK: [[META159]] = !DIDerivedType(tag: DW_TAG_rvalue_reference_type, baseType: [[META147]], size: 64, addressSpace: 1)
 // CHECK: [[META160]] = !{[[META161]]}
 // CHECK: [[META161]] = !DILocalVariable(arg: 1, scope: [[DBG144]], file: [[META4]], line: 226, type: [[META147]])
 // CHECK: [[META162]] = !DILocation(line: 226, column: 64, scope: [[DBG144]])
@@ -1584,7 +1584,7 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[DBG563]] = distinct !DISubprogram(name: "Test_Func_Pointer", linkageName: "_Z17Test_Func_PointerPi", scope: [[META4]], file: [[META4]], line: 782, type: [[META564:![0-9]+]], scopeLine: 782, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: [[META0]], retainedNodes: [[META567:![0-9]+]])
 // CHECK: [[META564]] = !DISubroutineType(types: [[META565:![0-9]+]])
 // CHECK: [[META565]] = !{null, [[META566:![0-9]+]]}
-// CHECK: [[META566]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META14]], size: 64)
+// CHECK: [[META566]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META14]], size: 64, addressSpace: 1)
 // CHECK: [[META567]] = !{[[META568]]}
 // CHECK: [[META568]] = !DILocalVariable(arg: 1, scope: [[DBG563]], file: [[META4]], line: 782, type: [[META566]])
 // CHECK: [[META569]] = !DILocation(line: 782, column: 44, scope: [[DBG563]])
@@ -1597,7 +1597,7 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[DBG576]] = distinct !DISubprogram(name: "Test_Func_Reference", linkageName: "_Z19Test_Func_ReferenceRi", scope: [[META4]], file: [[META4]], line: 808, type: [[META577:![0-9]+]], scopeLine: 808, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: [[META0]], retainedNodes: [[META580:![0-9]+]])
 // CHECK: [[META577]] = !DISubroutineType(types: [[META578:![0-9]+]])
 // CHECK: [[META578]] = !{null, [[META579:![0-9]+]]}
-// CHECK: [[META579]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META14]], size: 64)
+// CHECK: [[META579]] = !DIDerivedType(tag: DW_TAG_reference_type, baseType: [[META14]], size: 64, addressSpace: 1)
 // CHECK: [[META580]] = !{[[META581]]}
 // CHECK: [[META581]] = !DILocalVariable(arg: 1, scope: [[DBG576]], file: [[META4]], line: 808, type: [[META579]])
 // CHECK: [[META582]] = !DILocation(line: 808, column: 46, scope: [[DBG576]])
@@ -1629,7 +1629,7 @@ __global__ void Test_Kern_ParamRegLimitUnexpandedStruct(int64_t, int64_t, int64_
 // CHECK: [[META608]] = !{[[META609:![0-9]+]], [[META610:![0-9]+]]}
 // CHECK: [[META609]] = !DIDerivedType(tag: DW_TAG_member, name: "Element0", scope: [[META607]], file: [[META4]], line: 77, baseType: [[META566]], size: 64)
 // CHECK: [[META610]] = !DIDerivedType(tag: DW_TAG_member, name: "Element1", scope: [[META607]], file: [[META4]], line: 78, baseType: [[META611:![0-9]+]], size: 64, offset: 64)
-// CHECK: [[META611]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META612:![0-9]+]], size: 64)
+// CHECK: [[META611]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META612:![0-9]+]], size: 64, addressSpace: 1)
 // CHECK: [[META612]] = !DIBasicType(name: "float", size: 32, encoding: DW_ATE_float)
 // CHECK: [[META613]] = !{[[META614]]}
 // CHECK: [[META614]] = !DILocalVariable(arg: 1, scope: [[DBG604]], file: [[META4]], line: 859, type: [[META607]])

clang/test/CodeGenHIP/debug-info-diop-in-diexpression_ir.hip

@@ -125,7 +125,7 @@ __device__ void func1(pair p) {}
 // DIEXPRESSION-IR: [[META24]] = !DILocalVariable(name: "KernelVarA", scope: [[META19]], file: [[META7]], line: 50, type: [[META8]])
 // DIEXPRESSION-IR: [[META25]] = !DILocalVariable(name: "KernelVarB", scope: [[META19]], file: [[META7]], line: 51, type: [[META8]])
 // DIEXPRESSION-IR: [[META26]] = !DILocalVariable(name: "KernelVarSharedAPointer", scope: [[META19]], file: [[META7]], line: 53, type: [[META27:![0-9]+]])
-// DIEXPRESSION-IR: [[META27]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META8]], size: 64)
+// DIEXPRESSION-IR: [[META27]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: [[META8]], size: 64, addressSpace: 1)
 // DIEXPRESSION-IR: [[META28]] = !DILocalVariable(name: "KernelVarSharedBPointer", scope: [[META19]], file: [[META7]], line: 54, type: [[META27]])
 // DIEXPRESSION-IR: [[META29]] = !DILocalVariable(name: "KernelVarAPointer", scope: [[META19]], file: [[META7]], line: 55, type: [[META27]])
 // DIEXPRESSION-IR: [[META30]] = !DILocalVariable(name: "KernelVarBPointer", scope: [[META19]], file: [[META7]], line: 56, type: [[META27]])

clang/test/CodeGenHIP/debug-info-memory-space.hip

@@ -8,7 +8,7 @@
 
 // CHECK-DAG: !DILocalVariable(name: "FuncVarSharedPointer", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DEVICE_PTR:[0-9]+]])
 // CHECK-DAG: !DILocalVariable(name: "FuncVarPointer", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DEVICE_PTR:[0-9]+]])
-// CHECK-DAG: ![[DEVICE_PTR]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}})
+// CHECK-DAG: ![[DEVICE_PTR]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 1)
 
 #define __device__ __attribute__((device))
 #define __shared__ __attribute__((shared))

clang/test/CodeGenOpenCL/amdgpu-debug-info-pointer-address-space.cl

@@ -4,9 +4,9 @@
 
 // CHECK-DAG: ![[DWARF_MEMORY_SPACE_GLOBAL:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, memorySpace: DW_MSPACE_LLVM_global)
 // CHECK-DAG: ![[DWARF_MEMORY_SPACE_CONSTANT:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, memorySpace: DW_MSPACE_LLVM_constant)
-// CHECK-DAG: ![[DWARF_MEMORY_SPACE_LOCAL:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 2, memorySpace: DW_MSPACE_LLVM_group)
-// CHECK-DAG: ![[DWARF_MEMORY_SPACE_PRIVATE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 1, memorySpace: DW_MSPACE_LLVM_private)
-// CHECK-DAG: ![[DWARF_MEMORY_SPACE_NONE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}})
+// CHECK-DAG: ![[DWARF_MEMORY_SPACE_LOCAL:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 3, memorySpace: DW_MSPACE_LLVM_group)
+// CHECK-DAG: ![[DWARF_MEMORY_SPACE_PRIVATE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 5, memorySpace: DW_MSPACE_LLVM_private)
+// CHECK-DAG: ![[DWARF_MEMORY_SPACE_NONE:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, addressSpace: 1)
 
 // CHECK-DAG: distinct !DIGlobalVariable(name: "FileVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: ![[DWARF_MEMORY_SPACE_GLOBAL]], isLocal: false, isDefinition: true, memorySpace: DW_MSPACE_LLVM_global)
 global int *FileVar0;