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
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
