diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h --- a/clang/lib/Basic/Targets/NVPTX.h +++ b/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 diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -4664,15 +4664,7 @@ SmallVector Expr; unsigned AddressSpace = - CGM.getContext().getTargetAddressSpace(D->getType()); - if (CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) { - if (D->hasAttr()) - AddressSpace = - CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared); - else if (D->hasAttr()) - AddressSpace = - CGM.getContext().getTargetAddressSpace(LangAS::cuda_constant); - } + CGM.getContext().getTargetAddressSpace(CGM.GetGlobalVarAddressSpace(D)); AppendAddressSpaceXDeref(AddressSpace, Expr); GVE = DBuilder.createGlobalVariableExpression( diff --git a/clang/test/CodeGenHIP/debug-info-address-class.hip b/clang/test/CodeGenHIP/debug-info-address-class.hip new file mode 100644 --- /dev/null +++ b/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; +}