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
@@ -4667,15 +4667,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(
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,37 @@
+// 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]+}})