diff --git a/gdb/NEWS b/gdb/NEWS
index 57f8829d24f..74b0ba43986 100644
--- a/gdb/NEWS
+++ b/gdb/NEWS
@@ -116,6 +116,9 @@
   intent to remove it in a future release.
   The s390 64-bit target (s390x-*) remains supported.
 
+* GDB now recognizes DWARF DW_AT_address_space attribute for pointer,
+  reference, and rvalue reference types.
+
 * Configure changes
 
 ** --with-babeltrace has been removed.  The babeltrace library was
diff --git a/gdb/amdgpu-tdep.c b/gdb/amdgpu-tdep.c
index cfeddde7c4d..f4e4f5ee4c5 100644
--- a/gdb/amdgpu-tdep.c
+++ b/gdb/amdgpu-tdep.c
@@ -1598,6 +1598,71 @@ amdgpu_segment_address_from_core_address (CORE_ADDR addr)
   return addr & significant_bits;
 }
 
+/* Convert a given address in the given address space to an address in
+   another address space.  This is essentially a wrapper for a dbgapi
+   function.  PTID and SIMD_LANE are the contexts for which the
+   conversion should be made.  PTID can be NULL_PTID when the address
+   does not necessarily depend on wave context.  */
+
+static std::optional<addr_range>
+amdgpu_convert_address (gdbarch *gdbarch, ptid_t ptid, int simd_lane,
+			CORE_ADDR from_address,
+			arch_addr_space_id from_aspace_id,
+			arch_addr_space_id to_aspace_id)
+{
+  /* Argument should be already stripped off its address space id.  */
+  gdb_assert (amdgpu_address_space_id_from_core_address (from_address)
+	      == ARCH_ADDR_SPACE_ID_DEFAULT);
+
+  /* We should never get a valid ptid here that is not a gpu
+     thread.  */
+  gdb_assert (ptid == null_ptid || ptid_is_gpu (ptid));
+
+  const amd_dbgapi_wave_id_t wave_id
+    = (ptid != null_ptid)
+    ? get_amd_dbgapi_wave_id (ptid)
+    : AMD_DBGAPI_WAVE_NONE;
+
+  amd_dbgapi_architecture_id_t architecture_id;
+  if (ptid == null_ptid
+      && (amd_dbgapi_get_architecture (gdbarch_bfd_arch_info (gdbarch)->mach,
+				       &architecture_id)
+	  != AMD_DBGAPI_STATUS_SUCCESS))
+      error (_("amd_dbgapi_get_architecture failed"));
+  else if (ptid != null_ptid
+	   && (amd_dbgapi_wave_get_info (wave_id,
+					 AMD_DBGAPI_WAVE_INFO_ARCHITECTURE,
+					 sizeof (architecture_id),
+					 &architecture_id)
+	       != AMD_DBGAPI_STATUS_SUCCESS))
+    error (_("amd_dbgapi_wave_get_info failed"));
+
+  /* Get a dbgapi address space id for the original address space.  */
+  amd_dbgapi_address_space_id_t dbgapi_from_addr_space_id;
+  if (amd_dbgapi_dwarf_address_space_to_address_space
+	(architecture_id, from_aspace_id, &dbgapi_from_addr_space_id)
+      != AMD_DBGAPI_STATUS_SUCCESS)
+    error (_("amd_dbgapi_dwarf_address_space_to_address_space failed"));
+
+  /* Get the dbgapi address space id for the default address space.  */
+  amd_dbgapi_address_space_id_t dbgapi_to_addr_space_id;
+  if (amd_dbgapi_dwarf_address_space_to_address_space
+	(architecture_id, to_aspace_id, &dbgapi_to_addr_space_id)
+      != AMD_DBGAPI_STATUS_SUCCESS)
+    error (_("amd_dbgapi_dwarf_address_space_to_address_space failed"));
+
+  CORE_ADDR converted_address;
+  size_t contiguous_size;
+
+  if (amd_dbgapi_convert_address_space
+	(wave_id, simd_lane, dbgapi_from_addr_space_id, from_address,
+	 dbgapi_to_addr_space_id, &converted_address, &contiguous_size)
+      != AMD_DBGAPI_STATUS_SUCCESS)
+    return std::nullopt;
+
+  return addr_range {converted_address, contiguous_size};
+}
+
 /* Address class to address space mapping.
 
    TODO: This is just a quick fix to make the address class hand
@@ -1681,15 +1746,16 @@ amdgpu_address_class_id_to_name (struct gdbarch *gdbarch,
 /* Form a core address from a TYPE pointer type information and BUF
    pointer value buffer.
 
-   TODO: The address class information belongs to a type, which means
-	 that an address class is a language based concept, so either
-	 the language should define the address class numbers and
-	 names or some kind of a language enumeration needs to be
-	 passed in.
-
-	 At the moment, we assume that the language is OpenCL and we
-	 apply 1-1 mapping between its address classes and the AMDGPU
-	 address spaces.  */
+   TODO: TYPE may contain address space (i.e. DW_AT_address_space)
+   information.  In that case, we use it.  If empty, it is possible
+   that TYPE has address class (i.e. DW_AT_address_class) information.
+   OpenCL currently encodes source language address space qualifiers
+   in DW_AT_address_class.  Therefore we cannot ignore address class
+   and we handle them as "best effort".  Based on the direction DWARF
+   6 is taking, it is expected that OpenCL would stop using
+   DW_AT_address_class and use a new attribute instead.  When that
+   happens, all address-class related code, including gdbarch methods,
+   can be removed.  We just need address space.  */
 static CORE_ADDR
 amdgpu_pointer_to_address (struct gdbarch *gdbarch,
 			   struct type *type, const gdb_byte *buf)
@@ -1697,24 +1763,76 @@ amdgpu_pointer_to_address (struct gdbarch *gdbarch,
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
   CORE_ADDR address
     = extract_unsigned_integer (buf, type->length (), byte_order);
-  unsigned int aclass_id = type->address_class ();
-  unsigned int address_class
-    = amdgpu_type_flags_aclass_to_addr_class (aclass_id);
-
-  /* Address might be in a converted format already, so even if the
-     class is global, the address might have the address space part
-     in it.  This happens in cases like 'p &local_array'."  */
-  if (address_class == DWARF_GLOBAL_ADDR_CLASS)
-    return address;
 
-  /* In the current implementation, we shouldn't have a case where we
-     have both type address class information as well as address
-     space information in a core address.  */
+  /* No address space information should be encoded in the pointer
+     value.  We use the pointer type for that purpose.  */
   gdb_assert (!amdgpu_address_space_id_from_core_address (address));
 
+  arch_addr_space_id address_space = type->address_space ();
+  if (address_space == 0)
+    {
+      /* Try address class.  At the moment, we assume 1-1 mapping
+	 between address classes and AMDGPU address spaces.  */
+      unsigned int aclass_id = type->address_class ();
+      address_space
+	= amdgpu_type_flags_aclass_to_addr_class (aclass_id);
+    }
+
+  return amdgpu_segment_address_to_core_address (address_space, address);
+}
+
+/* Store ADDRESS, which may have an address space id encoded in it, to
+   an address value of type TYPE, which also has address space
+   information.  */
+
+static void
+amdgpu_address_to_pointer (gdbarch *gdbarch, type *type,
+			   gdb_byte *buffer, CORE_ADDR address)
+{
+  arch_addr_space_id addr_aspace
+    = amdgpu_address_space_id_from_core_address (address);
+  /* There is 1-1 mapping between DWARF address spaces and GPU address
+     spaces.  */
+  arch_addr_space_id type_aspace = type->address_space ();
+
+  gdb_assert (addr_aspace == type_aspace);
+
   address = amdgpu_segment_address_from_core_address (address);
+  unsigned_address_to_pointer (gdbarch, type, buffer, address);
+}
 
-  return amdgpu_segment_address_to_core_address (address_class, address);
+/* Convert ADDRESS from one address space to another.  */
+static CORE_ADDR
+amdgpu_pointer_to_pointer (gdbarch *gdbarch, type *from_type,
+			   CORE_ADDR address, type *to_type)
+{
+  arch_addr_space_id from_type_aspace = from_type->address_space ();
+  arch_addr_space_id to_type_aspace = to_type->address_space ();
+
+  /* Strip address space info from the address.  Initialize
+     FROM_ADDRESS to 0, although it is written by
+     'address_to_pointer', because FROM_TYPE may be shorter than size
+     of CORE_ADDR.  */
+  CORE_ADDR from_address = 0;
+  amdgpu_address_to_pointer (gdbarch, from_type,
+			     (gdb_byte *) &from_address, address);
+
+  /* Convert.  */
+  ptid_t ptid = inferior_thread ()->ptid;
+  gdb_assert (ptid_is_gpu (ptid));
+  int lane = inferior_thread ()->current_simd_lane ();
+
+  std::optional<addr_range> range
+    = amdgpu_convert_address (gdbarch, ptid, lane, from_address,
+			      from_type_aspace, to_type_aspace);
+  if (!range.has_value ())
+    error (_("Cannot convert pointer-to-'%s' to a pointer-to-'%s'"),
+	   gdbarch_address_space_id_to_name (gdbarch, from_type_aspace),
+	   gdbarch_address_space_id_to_name (gdbarch, to_type_aspace));
+
+  /* Add address space info.  */
+  return amdgpu_pointer_to_address (gdbarch, to_type,
+				    (gdb_byte *) &(range.value ().addr));
 }
 
 static CORE_ADDR
@@ -1750,6 +1868,47 @@ amdgpu_address_spaces (struct gdbarch *gdbarch)
   return tdep->address_spaces;
 }
 
+/* Implementation of the 'address_space_pointer_size' gdbarch method.  */
+
+static unsigned int
+amdgpu_address_space_pointer_size (gdbarch *gdbarch,
+				   arch_addr_space_id aspace)
+{
+  amd_dbgapi_architecture_id_t architecture_id;
+  if (inferior_ptid != null_ptid
+      && ptid_is_gpu (inferior_thread ()->ptid))
+    {
+      const amd_dbgapi_wave_id_t wave_id
+	= get_amd_dbgapi_wave_id (inferior_thread ()->ptid);
+      if (amd_dbgapi_wave_get_info (wave_id,
+				    AMD_DBGAPI_WAVE_INFO_ARCHITECTURE,
+				    sizeof (architecture_id),
+				    &architecture_id)
+	  != AMD_DBGAPI_STATUS_SUCCESS)
+	error (_("amd_dbgapi_get_architecture failed"));
+    }
+  else if (amd_dbgapi_get_architecture (gdbarch_bfd_arch_info (gdbarch)->mach,
+					&architecture_id)
+	   != AMD_DBGAPI_STATUS_SUCCESS)
+    error (_("amd_dbgapi_get_architecture failed"));
+
+  /* Get a dbgapi address space id for the original address space.  */
+  amd_dbgapi_address_space_id_t dbgapi_aspace_id;
+  if (amd_dbgapi_dwarf_address_space_to_address_space
+	(architecture_id, aspace, &dbgapi_aspace_id)
+      != AMD_DBGAPI_STATUS_SUCCESS)
+    error (_("amd_dbgapi_dwarf_address_space_to_address_space failed"));
+
+  amd_dbgapi_size_t size;
+  if (amd_dbgapi_address_space_get_info
+	(dbgapi_aspace_id, AMD_DBGAPI_ADDRESS_SPACE_INFO_ADDRESS_SIZE,
+	 sizeof (size), &size)
+      != AMD_DBGAPI_STATUS_SUCCESS)
+    error (_("amd_dbgapi_address_space_get_info failed"));
+
+  return (unsigned int) size / TARGET_CHAR_BIT;
+}
+
 static location_scope
 amdgpu_address_scope (struct gdbarch *gdbarch, ptid_t ptid, CORE_ADDR address)
 {
@@ -1845,67 +2004,26 @@ amdgpu_get_watchable_aliases (struct gdbarch *gdbarch,
   if (size == 0)
     error (_("Watchpoint length must be non-zero number."));
 
-  /* We should never get a valid ptid here that is not a gpu
-     thread.  */
-  gdb_assert (ptid == null_ptid || ptid_is_gpu (ptid));
-
   std::vector<addr_range> aliases;
-  amd_dbgapi_architecture_id_t architecture_id;
-  /* If we have a wave, use the wave's arch, otherwise, use the arch matching
-     gdbarch.  */
-  if (ptid != null_ptid)
-    {
-      amd_dbgapi_wave_id_t wave_id = get_amd_dbgapi_wave_id (ptid);
-      if (amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_ARCHITECTURE,
-				    sizeof (architecture_id), &architecture_id)
-	  != AMD_DBGAPI_STATUS_SUCCESS)
-	error (_("amd_dbgapi_wave_get_info failed"));
-    }
-  else
-    {
-      if (amd_dbgapi_get_architecture (gdbarch_bfd_arch_info (gdbarch)->mach,
-				       &architecture_id)
-	  != AMD_DBGAPI_STATUS_SUCCESS)
-	error (_("amd_dbgapi_get_architecture failed"));
-    }
-
-  /* Get a dbgapi address space id for the original address space.  */
-  amd_dbgapi_address_space_id_t dbgapi_from_addr_space_id;
-  if (amd_dbgapi_dwarf_address_space_to_address_space
-	(architecture_id, (uint64_t) addr_space_id,
-	 &dbgapi_from_addr_space_id)
-	!= AMD_DBGAPI_STATUS_SUCCESS)
-    error (_("amd_dbgapi_dwarf_address_space_to_address_space failed"));
-
-  /* Get the dbgapi address space id for the default address space.  */
-  amd_dbgapi_address_space_id_t dbgapi_to_addr_space_id;
-  gdb_assert (amd_dbgapi_dwarf_address_space_to_address_space
-	       (architecture_id, (uint64_t) ARCH_ADDR_SPACE_ID_DEFAULT,
-		&dbgapi_to_addr_space_id)
-	       == AMD_DBGAPI_STATUS_SUCCESS);
-
-  amd_dbgapi_wave_id_t wave_id = AMD_DBGAPI_WAVE_NONE;
-
-  if (ptid != null_ptid)
-    wave_id = get_amd_dbgapi_wave_id (ptid);
 
   /* Aliasing address range might not have the same layout.  Because
      of that, when ever there is a gap between the aliasing addresses,
      create a new range.  */
   while (true)
     {
-      amd_dbgapi_size_t converted_size;
-      amd_dbgapi_segment_address_t to_offset;
-
       /* Try to convert the address to an address in the
 	 default address space.  */
-      if (amd_dbgapi_convert_address_space
-	    (wave_id, simd_lane, dbgapi_from_addr_space_id,
-	     (amd_dbgapi_segment_address_t) addr,
-	     dbgapi_to_addr_space_id, &to_offset, &converted_size)
-	    != AMD_DBGAPI_STATUS_SUCCESS)
+      std::optional<addr_range> converted_range
+	= amdgpu_convert_address (gdbarch, ptid, simd_lane, addr,
+				  addr_space_id,
+				  ARCH_ADDR_SPACE_ID_DEFAULT);
+
+      if (!converted_range.has_value ())
 	return {};
 
+      size_t converted_size = converted_range.value ().size;
+      CORE_ADDR to_offset = converted_range.value ().addr;
+
       if (converted_size == 0)
 	return {};
 
@@ -2069,6 +2187,10 @@ amdgpu_gdbarch_init (struct gdbarch_info info, struct gdbarch_list *arches)
   set_gdbarch_dummy_id (gdbarch, amdgpu_dummy_id);
 
   set_gdbarch_pointer_to_address (gdbarch, amdgpu_pointer_to_address);
+  set_gdbarch_address_to_pointer (gdbarch, amdgpu_address_to_pointer);
+  set_gdbarch_pointer_to_pointer (gdbarch, amdgpu_pointer_to_pointer);
+  set_gdbarch_address_space_pointer_size
+    (gdbarch, amdgpu_address_space_pointer_size);
   set_gdbarch_address_class_dwarf_to_id
     (gdbarch, amdgpu_address_class_dwarf_to_id);
   set_gdbarch_address_class_id_to_name
diff --git a/gdb/arch-utils.c b/gdb/arch-utils.c
index a820128894b..bf0fd0b4550 100644
--- a/gdb/arch-utils.c
+++ b/gdb/arch-utils.c
@@ -1087,7 +1087,7 @@ default_read_core_file_mappings
 
 /* See arch-utils.h.  */
 std::optional<arch_addr_space_id>
-gdbarch_name_to_address_space_id (struct gdbarch *gdbarch, const char *name)
+gdbarch_address_space_name_to_id (struct gdbarch *gdbarch, const char *name)
 {
   if (!gdbarch_address_spaces_p (gdbarch))
     return {};
@@ -1148,9 +1148,19 @@ non-default address spaces."));
 
 /* See arch-utils.h.  */
 arch_addr_space_id
-default_dwarf_address_space_to_address_space_id (LONGEST dwarf_addr_space)
+default_address_space_dwarf_to_id (gdbarch *gdbarch,
+				   ULONGEST dwarf_addr_space)
 {
-  return (arch_addr_space_id) dwarf_addr_space;
+  if (gdbarch_address_spaces_p (gdbarch))
+    {
+      /* By default, assume 1-1 mapping.  */
+      for (const auto &address_space : gdbarch_address_spaces (gdbarch))
+	if (address_space.id == dwarf_addr_space)
+	  return (arch_addr_space_id) dwarf_addr_space;
+    }
+
+  error (_("DWARF address space id (%s) is not recognized by "
+	   "the architecture."), pulongest (dwarf_addr_space));
 }
 
 /* See arch-utils.h.  */
diff --git a/gdb/arch-utils.h b/gdb/arch-utils.h
index cb259a3cd7e..9b8ef7e02c4 100644
--- a/gdb/arch-utils.h
+++ b/gdb/arch-utils.h
@@ -414,7 +414,7 @@ extern std::optional<CORE_ADDR> default_get_shadow_stack_pointer
   /* Architecture address space handling.  */
 
 /* Convert address space name to an address space id.  */
-extern std::optional<arch_addr_space_id> gdbarch_name_to_address_space_id
+extern std::optional<arch_addr_space_id> gdbarch_address_space_name_to_id
   (struct gdbarch *gdbarch, const char *name);
 
 /* Convert address space id to an address space name.  */
@@ -422,9 +422,9 @@ extern const char *gdbarch_address_space_id_to_name
   (struct gdbarch *gdbarch, arch_addr_space_id addr_space_id);
 
 /* Default implementation of
-   gdbarch_dwarf_address_space_to_address_space_id.  */
-extern arch_addr_space_id default_dwarf_address_space_to_address_space_id
-  (LONGEST dwarf_addr_space);
+   gdbarch_address_space_dwarf_to_id.  */
+extern arch_addr_space_id default_address_space_dwarf_to_id
+(gdbarch *gdbarch, ULONGEST dwarf_addr_space);
 
 /* TODO: Following default address space hooks are a quick fix until a proper
 	 address space support is added and should not be pushed upstream.  */
diff --git a/gdb/dwarf2/expr.c b/gdb/dwarf2/expr.c
index 798215f59ed..a0150c54ff3 100644
--- a/gdb/dwarf2/expr.c
+++ b/gdb/dwarf2/expr.c
@@ -3681,7 +3681,7 @@ dwarf_expr_context::execute_llvm_stack_op (dwarf_llvm_user op,
 
 	  dwarf_require_integral (aspace_value->get_type ());
 	  arch_addr_space_id address_space
-	    = gdbarch_dwarf_address_space_to_address_space_id
+	    = gdbarch_address_space_dwarf_to_id
 		(arch, aspace_value->to_long ());
 	  CORE_ADDR address = address_value->to_long ();
 	  address
@@ -3834,7 +3834,7 @@ dwarf_expr_context::execute_llvm_stack_op (dwarf_llvm_user op,
 
 	  dwarf_require_integral (aspace_value->get_type ());
 	  arch_addr_space_id address_space
-	    = gdbarch_dwarf_address_space_to_address_space_id
+	    = gdbarch_address_space_dwarf_to_id
 		(arch, aspace_value->to_long ());
 	  memory->set_address_space (address_space);
 	  result_entry = memory;
diff --git a/gdb/dwarf2/read.c b/gdb/dwarf2/read.c
index 9461c3c5802..c59fc6ee697 100644
--- a/gdb/dwarf2/read.c
+++ b/gdb/dwarf2/read.c
@@ -12001,6 +12001,30 @@ is_recursive_pointer (die_info *die, dwarf2_cu *cu)
   return true;
 }
 
+/* Helper function to read DW_AT_address_space attribute.  */
+
+static ULONGEST
+get_address_space (die_info *die, dwarf2_cu *cu)
+{
+  ULONGEST addr_space;
+  if (attribute *attr_aspace = dwarf2_attr (die, DW_AT_LLVM_address_space, cu);
+      attr_aspace != nullptr)
+    addr_space = (attr_aspace->unsigned_constant ()
+		  .value_or (DW_ASPACE_default));
+  else
+    addr_space = DW_ASPACE_default;
+
+  if (addr_space > type_instance_flags::ADDRESS_SPACE_MAX)
+    {
+      warning (_("address space value (%s) obtained from DWARF exceeds "
+		 "max encodable value; using default address space"),
+	       pulongest (addr_space));
+      addr_space = DW_ASPACE_default;
+    }
+
+  return addr_space;
+}
+
 /* Extract all information from a DW_TAG_pointer_type DIE and add to
    the user defined type vector.  */
 
@@ -12046,12 +12070,13 @@ read_tag_pointer_type (struct die_info *die, struct dwarf2_cu *cu)
   else
     addr_class = DW_ADDR_none;
 
+  ULONGEST addr_space = get_address_space (die, cu);
   ULONGEST alignment = get_alignment (cu, die);
 
-  /* If the pointer size, alignment, or address class is different
+  /* If the pointer size, alignment, or address class/space is different
      than the default, create a type variant marked as such and set
      the length accordingly.  */
-  if (addr_class != DW_ADDR_none)
+  if (addr_class != DW_ADDR_none || addr_space != DW_ASPACE_default)
     {
       if (gdbarch_address_class_dwarf_to_id_p (gdbarch))
 	{
@@ -12060,10 +12085,26 @@ read_tag_pointer_type (struct die_info *die, struct dwarf2_cu *cu)
 						 addr_class);
 	  type = make_type_with_address_class (type, aclass);
 	}
-      else
+      else if (addr_class != DW_ADDR_none)
 	{
 	  /* Should we also complain about unhandled address classes?  */
 	}
+
+      arch_addr_space_id aspace
+	= gdbarch_address_space_dwarf_to_id (gdbarch, addr_space);
+      type = make_type_with_address_space (type, aspace);
+
+      /* We expect the type length information coming from DWARF to
+	 match what the architecture says.  If there is a mismatch,
+	 issue a complaint and continue with arch's decision.  */
+      if (type->length () != byte_size)
+	{
+	  complaint (_("length of pointer with address space (%s) and "
+		       "address class (%s) is %s in DWARF but arch says %s"),
+		     pulongest (addr_space), pulongest (addr_class),
+		     pulongest (byte_size), pulongest (type->length ()));
+	  byte_size = type->length ();
+	}
     }
   else if (type->length () != byte_size)
     complaint (_("invalid pointer size %s"), pulongest (byte_size));
@@ -12129,6 +12170,7 @@ read_tag_reference_type (struct die_info *die, struct dwarf2_cu *cu,
 			  enum type_code refcode)
 {
   unit_head *cu_header = &cu->header;
+  gdbarch *gdbarch = cu->per_objfile->objfile->arch ();
   struct type *type, *target_type;
   struct attribute *attr;
 
@@ -12142,6 +12184,15 @@ read_tag_reference_type (struct die_info *die, struct dwarf2_cu *cu,
     return type;
 
   type = lookup_reference_type (target_type, refcode);
+
+  ULONGEST addr_space = get_address_space (die, cu);
+  if (addr_space != DW_ASPACE_default)
+    {
+      arch_addr_space_id aspace
+	= gdbarch_address_space_dwarf_to_id (gdbarch, addr_space);
+      type = make_type_with_address_space (type, aspace);
+    }
+
   attr = dwarf2_attr (die, DW_AT_byte_size, cu);
   if (attr != nullptr)
     type->set_length (attr->unsigned_constant ()
diff --git a/gdb/eval.c b/gdb/eval.c
index 64f06388e18..26aa34bb951 100644
--- a/gdb/eval.c
+++ b/gdb/eval.c
@@ -2564,7 +2564,7 @@ aspace_operation::evaluate (struct type *expect_type,
 
   const std::string &name = std::get<1> (m_storage);
   std::optional<arch_addr_space_id> address_space_id
-    = gdbarch_name_to_address_space_id (exp->gdbarch, name.c_str ());
+    = gdbarch_address_space_name_to_id (exp->gdbarch, name.c_str ());
 
   if (!address_space_id.has_value ())
     error (_("Address space %s not recognised by the architecture"),
diff --git a/gdb/gdb-gdb.py.in b/gdb/gdb-gdb.py.in
index 7e39d27366e..b0080e47513 100644
--- a/gdb/gdb-gdb.py.in
+++ b/gdb/gdb-gdb.py.in
@@ -38,6 +38,8 @@ class StructTypeInstanceFlagsPrettyPrinter:
             fields.append("DATA_SPACE")
         if self.val["address_class"] != 0:
             fields.append("ADDRESS_CLASS(%d)" % self.val["address_class"])
+        if self.val["address_space"] != 0:
+            fields.append("ADDRESS_SPACE(%d)" % self.val["address_space"])
         if self.val["is_nottext"]:
             fields.append("NOTTEXT")
         if self.val["is_restrict"]:
diff --git a/gdb/gdbarch-gen.c b/gdb/gdbarch-gen.c
index 41656f17534..c50d7da635c 100644
--- a/gdb/gdbarch-gen.c
+++ b/gdb/gdbarch-gen.c
@@ -110,12 +110,14 @@ struct gdbarch
   gdbarch_dwarf2_reg_piece_offset_ftype *dwarf2_reg_piece_offset = default_dwarf2_reg_piece_offset;
   gdbarch_pointer_to_address_ftype *pointer_to_address = unsigned_pointer_to_address;
   gdbarch_address_to_pointer_ftype *address_to_pointer = unsigned_address_to_pointer;
+  gdbarch_pointer_to_pointer_ftype *pointer_to_pointer = nullptr;
   gdbarch_integer_to_address_ftype *integer_to_address = nullptr;
   gdbarch_address_spaces_ftype *address_spaces = nullptr;
   gdbarch_address_space_id_from_core_address_ftype *address_space_id_from_core_address = nullptr;
   gdbarch_segment_address_from_core_address_ftype *segment_address_from_core_address = nullptr;
   gdbarch_segment_address_to_core_address_ftype *segment_address_to_core_address = nullptr;
-  gdbarch_dwarf_address_space_to_address_space_id_ftype *dwarf_address_space_to_address_space_id = nullptr;
+  gdbarch_address_space_dwarf_to_id_ftype *address_space_dwarf_to_id = nullptr;
+  gdbarch_address_space_pointer_size_ftype *address_space_pointer_size = nullptr;
   gdbarch_address_scope_ftype *address_scope = nullptr;
   gdbarch_get_watchable_aliases_ftype *get_watchable_aliases = default_get_watchable_aliases;
   gdbarch_return_value_ftype *return_value = nullptr;
@@ -373,6 +375,7 @@ verify_gdbarch (struct gdbarch *gdbarch)
   /* Skip verify of dwarf2_reg_piece_offset, invalid_p == 0.  */
   /* Skip verify of pointer_to_address, invalid_p == 0.  */
   /* Skip verify of address_to_pointer, invalid_p == 0.  */
+  /* Skip verify of pointer_to_pointer, has predicate.  */
   /* Skip verify of integer_to_address, has predicate.  */
   /* Skip verify of address_spaces, has predicate.  */
   if (gdbarch->address_space_id_from_core_address == nullptr)
@@ -384,9 +387,10 @@ verify_gdbarch (struct gdbarch *gdbarch)
   if (gdbarch->segment_address_to_core_address == nullptr)
     gdbarch->segment_address_to_core_address = default_segment_address_to_core_address;
   /* Skip verify of segment_address_to_core_address, invalid_p == 0.  */
-  if (gdbarch->dwarf_address_space_to_address_space_id == nullptr)
-    gdbarch->dwarf_address_space_to_address_space_id = default_dwarf_address_space_to_address_space_id;
-  /* Skip verify of dwarf_address_space_to_address_space_id, invalid_p == 0.  */
+  if (gdbarch->address_space_dwarf_to_id == nullptr)
+    gdbarch->address_space_dwarf_to_id = default_address_space_dwarf_to_id;
+  /* Skip verify of address_space_dwarf_to_id, invalid_p == 0.  */
+  /* Skip verify of address_space_pointer_size, has predicate.  */
   if (gdbarch->address_scope == nullptr)
     gdbarch->address_scope = default_address_scope;
   /* Skip verify of address_scope, invalid_p == 0.  */
@@ -803,6 +807,12 @@ gdbarch_dump (struct gdbarch *gdbarch, struct ui_file *file)
   gdb_printf (file,
 	      "gdbarch_dump: address_to_pointer = <%s>\n",
 	      host_address_to_string (gdbarch->address_to_pointer));
+  gdb_printf (file,
+	      "gdbarch_dump: gdbarch_pointer_to_pointer_p() = %d\n",
+	      gdbarch_pointer_to_pointer_p (gdbarch));
+  gdb_printf (file,
+	      "gdbarch_dump: pointer_to_pointer = <%s>\n",
+	      host_address_to_string (gdbarch->pointer_to_pointer));
   gdb_printf (file,
 	      "gdbarch_dump: gdbarch_integer_to_address_p() = %d\n",
 	      gdbarch_integer_to_address_p (gdbarch));
@@ -825,8 +835,14 @@ gdbarch_dump (struct gdbarch *gdbarch, struct ui_file *file)
 	      "gdbarch_dump: segment_address_to_core_address = <%s>\n",
 	      host_address_to_string (gdbarch->segment_address_to_core_address));
   gdb_printf (file,
-	      "gdbarch_dump: dwarf_address_space_to_address_space_id = <%s>\n",
-	      host_address_to_string (gdbarch->dwarf_address_space_to_address_space_id));
+	      "gdbarch_dump: address_space_dwarf_to_id = <%s>\n",
+	      host_address_to_string (gdbarch->address_space_dwarf_to_id));
+  gdb_printf (file,
+	      "gdbarch_dump: gdbarch_address_space_pointer_size_p() = %d\n",
+	      gdbarch_address_space_pointer_size_p (gdbarch));
+  gdb_printf (file,
+	      "gdbarch_dump: address_space_pointer_size = <%s>\n",
+	      host_address_to_string (gdbarch->address_space_pointer_size));
   gdb_printf (file,
 	      "gdbarch_dump: address_scope = <%s>\n",
 	      host_address_to_string (gdbarch->address_scope));
@@ -2616,6 +2632,30 @@ set_gdbarch_address_to_pointer (struct gdbarch *gdbarch,
   gdbarch->address_to_pointer = address_to_pointer;
 }
 
+bool
+gdbarch_pointer_to_pointer_p (struct gdbarch *gdbarch)
+{
+  gdb_assert (gdbarch != nullptr);
+  return gdbarch->pointer_to_pointer != nullptr;
+}
+
+CORE_ADDR
+gdbarch_pointer_to_pointer (struct gdbarch *gdbarch, type *from_type, CORE_ADDR address, type *to_type)
+{
+  gdb_assert (gdbarch != nullptr);
+  gdb_assert (gdbarch->pointer_to_pointer != nullptr);
+  if (gdbarch_debug >= 2)
+    gdb_printf (gdb_stdlog, "gdbarch_pointer_to_pointer called\n");
+  return gdbarch->pointer_to_pointer (gdbarch, from_type, address, to_type);
+}
+
+void
+set_gdbarch_pointer_to_pointer (struct gdbarch *gdbarch,
+				gdbarch_pointer_to_pointer_ftype pointer_to_pointer)
+{
+  gdbarch->pointer_to_pointer = pointer_to_pointer;
+}
+
 bool
 gdbarch_integer_to_address_p (struct gdbarch *gdbarch)
 {
@@ -2716,20 +2756,44 @@ set_gdbarch_segment_address_to_core_address (struct gdbarch *gdbarch,
 }
 
 arch_addr_space_id
-gdbarch_dwarf_address_space_to_address_space_id (struct gdbarch *gdbarch, LONGEST dwarf_addr_space)
+gdbarch_address_space_dwarf_to_id (struct gdbarch *gdbarch, ULONGEST dwarf_addr_space)
+{
+  gdb_assert (gdbarch != nullptr);
+  gdb_assert (gdbarch->address_space_dwarf_to_id != nullptr);
+  if (gdbarch_debug >= 2)
+    gdb_printf (gdb_stdlog, "gdbarch_address_space_dwarf_to_id called\n");
+  return gdbarch->address_space_dwarf_to_id (gdbarch, dwarf_addr_space);
+}
+
+void
+set_gdbarch_address_space_dwarf_to_id (struct gdbarch *gdbarch,
+				       gdbarch_address_space_dwarf_to_id_ftype address_space_dwarf_to_id)
+{
+  gdbarch->address_space_dwarf_to_id = address_space_dwarf_to_id;
+}
+
+bool
+gdbarch_address_space_pointer_size_p (struct gdbarch *gdbarch)
+{
+  gdb_assert (gdbarch != nullptr);
+  return gdbarch->address_space_pointer_size != nullptr;
+}
+
+unsigned int
+gdbarch_address_space_pointer_size (struct gdbarch *gdbarch, arch_addr_space_id aspace)
 {
   gdb_assert (gdbarch != nullptr);
-  gdb_assert (gdbarch->dwarf_address_space_to_address_space_id != nullptr);
+  gdb_assert (gdbarch->address_space_pointer_size != nullptr);
   if (gdbarch_debug >= 2)
-    gdb_printf (gdb_stdlog, "gdbarch_dwarf_address_space_to_address_space_id called\n");
-  return gdbarch->dwarf_address_space_to_address_space_id (dwarf_addr_space);
+    gdb_printf (gdb_stdlog, "gdbarch_address_space_pointer_size called\n");
+  return gdbarch->address_space_pointer_size (gdbarch, aspace);
 }
 
 void
-set_gdbarch_dwarf_address_space_to_address_space_id (struct gdbarch *gdbarch,
-						     gdbarch_dwarf_address_space_to_address_space_id_ftype dwarf_address_space_to_address_space_id)
+set_gdbarch_address_space_pointer_size (struct gdbarch *gdbarch,
+					gdbarch_address_space_pointer_size_ftype address_space_pointer_size)
 {
-  gdbarch->dwarf_address_space_to_address_space_id = dwarf_address_space_to_address_space_id;
+  gdbarch->address_space_pointer_size = address_space_pointer_size;
 }
 
 location_scope
diff --git a/gdb/gdbarch-gen.h b/gdb/gdbarch-gen.h
index 539cb3db00a..dc02ef3849d 100644
--- a/gdb/gdbarch-gen.h
+++ b/gdb/gdbarch-gen.h
@@ -443,6 +443,15 @@ using gdbarch_address_to_pointer_ftype = void (struct gdbarch *gdbarch, struct t
 void gdbarch_address_to_pointer (struct gdbarch *gdbarch, struct type *type, gdb_byte *buf, CORE_ADDR addr);
 void set_gdbarch_address_to_pointer (struct gdbarch *gdbarch, gdbarch_address_to_pointer_ftype *address_to_pointer);
 
+/* Convert an address of type FROM_TYPE to an address of type TO_TYPE.
+   This is particularly useful to convert between address spaces. */
+
+bool gdbarch_pointer_to_pointer_p (struct gdbarch *gdbarch);
+
+using gdbarch_pointer_to_pointer_ftype = CORE_ADDR (struct gdbarch *gdbarch, type *from_type, CORE_ADDR address, type *to_type);
+CORE_ADDR gdbarch_pointer_to_pointer (struct gdbarch *gdbarch, type *from_type, CORE_ADDR address, type *to_type);
+void set_gdbarch_pointer_to_pointer (struct gdbarch *gdbarch, gdbarch_pointer_to_pointer_ftype *pointer_to_pointer);
+
 bool gdbarch_integer_to_address_p (struct gdbarch *gdbarch);
 
 using gdbarch_integer_to_address_ftype = CORE_ADDR (struct gdbarch *gdbarch, struct type *type, const gdb_byte *buf, arch_addr_space_id address_space_id);
@@ -481,13 +490,22 @@ using gdbarch_segment_address_to_core_address_ftype = CORE_ADDR (arch_addr_space
 CORE_ADDR gdbarch_segment_address_to_core_address (struct gdbarch *gdbarch, arch_addr_space_id address_space_id, CORE_ADDR address);
 void set_gdbarch_segment_address_to_core_address (struct gdbarch *gdbarch, gdbarch_segment_address_to_core_address_ftype *segment_address_to_core_address);
 
-/* Converts DWARF address space number to address space id.
-   TODO: This hook is a quick fix until a proper address space support
-   is added and should not be pushed upstream. */
+/* Given the DWARF identifier for an architecture-specific address space,
+   return the id of that address space. */
+
+using gdbarch_address_space_dwarf_to_id_ftype = arch_addr_space_id (struct gdbarch *gdbarch, ULONGEST dwarf_addr_space);
+arch_addr_space_id gdbarch_address_space_dwarf_to_id (struct gdbarch *gdbarch, ULONGEST dwarf_addr_space);
+void set_gdbarch_address_space_dwarf_to_id (struct gdbarch *gdbarch, gdbarch_address_space_dwarf_to_id_ftype *address_space_dwarf_to_id);
+
+/* Given an architecture-specific address space, return the size of a pointer
+   that points to that address space.  This size is not necessarily the same
+   as the size of a default pointer. */
+
+bool gdbarch_address_space_pointer_size_p (struct gdbarch *gdbarch);
 
-using gdbarch_dwarf_address_space_to_address_space_id_ftype = arch_addr_space_id (LONGEST dwarf_addr_space);
-arch_addr_space_id gdbarch_dwarf_address_space_to_address_space_id (struct gdbarch *gdbarch, LONGEST dwarf_addr_space);
-void set_gdbarch_dwarf_address_space_to_address_space_id (struct gdbarch *gdbarch, gdbarch_dwarf_address_space_to_address_space_id_ftype *dwarf_address_space_to_address_space_id);
+using gdbarch_address_space_pointer_size_ftype = unsigned int (struct gdbarch *gdbarch, arch_addr_space_id aspace);
+unsigned int gdbarch_address_space_pointer_size (struct gdbarch *gdbarch, arch_addr_space_id aspace);
+void set_gdbarch_address_space_pointer_size (struct gdbarch *gdbarch, gdbarch_address_space_pointer_size_ftype *address_space_pointer_size);
 
 /* Return the address's scope. */
 
diff --git a/gdb/gdbarch_components.py b/gdb/gdbarch_components.py
index d24d77896e4..87946adf569 100644
--- a/gdb/gdbarch_components.py
+++ b/gdb/gdbarch_components.py
@@ -854,6 +854,18 @@
     invalid=False,
 )
 
+Method(
+    comment="""
+Convert an address of type FROM_TYPE to an address of type TO_TYPE.
+This is particularly useful to convert between address spaces.
+""",
+    type="CORE_ADDR",
+    name="pointer_to_pointer",
+    params=[("type *", "from_type"), ("CORE_ADDR", "address"), ("type *", "to_type")],
+    predicate=True,
+    invalid=True,
+)
+
 Method(
     type="CORE_ADDR",
     name="integer_to_address",
@@ -915,19 +927,30 @@
     invalid=False,
 )
 
-Function(
+Method(
     comment="""
-Converts DWARF address space number to address space id.
-TODO: This hook is a quick fix until a proper address space support
-is added and should not be pushed upstream.
+Given the DWARF identifier for an architecture-specific address space,
+return the id of that address space.
 """,
     type="arch_addr_space_id",
-    name="dwarf_address_space_to_address_space_id",
-    params=[("LONGEST", "dwarf_addr_space")],
-    postdefault="default_dwarf_address_space_to_address_space_id",
+    name="address_space_dwarf_to_id",
+    params=[("ULONGEST", "dwarf_addr_space")],
+    postdefault="default_address_space_dwarf_to_id",
     invalid=False,
 )
 
+Method(
+    comment="""
+Given an architecture-specific address space, return the size of a pointer
+that points to that address space.  This size is not necessarily the same
+as the size of a default pointer.
+""",
+    type="unsigned int",
+    name="address_space_pointer_size",
+    params=[("arch_addr_space_id", "aspace")],
+    predicate=True,
+)
+
 Method(
     comment="""
 Return the address's scope.
diff --git a/gdb/gdbtypes.c b/gdb/gdbtypes.c
index f5c0177bff2..c62391ad27b 100644
--- a/gdb/gdbtypes.c
+++ b/gdb/gdbtypes.c
@@ -560,8 +560,15 @@ make_qualified_type (struct type *type, type_instance_flags new_flags,
   /* Now set the instance flags and return the new type.  */
   ntype->set_instance_flags (new_flags);
 
-  /* Set length of new type to that of the original type.  */
-  ntype->set_length (type->length ());
+  /* Set length of new type to that of the original type, except for
+     pointers to address spaces, which may have a different size.  */
+  if (ntype->address_space () != 0
+      && gdbarch_address_space_pointer_size_p (type->arch ()))
+    ntype->set_length
+      (gdbarch_address_space_pointer_size (type->arch (),
+					   ntype->address_space ()));
+  else
+    ntype->set_length (type->length ());
 
   return ntype;
 }
@@ -605,6 +612,18 @@ make_type_with_address_class (struct type *type,
   return make_qualified_type (type, new_flags, nullptr);
 }
 
+/* Make a variant of TYPE with ASPACE as the type's
+   DW_AT_address_space attribute.  */
+type *
+make_type_with_address_space (type *type, arch_addr_space_id aspace)
+{
+  gdb_assert (aspace <= type_instance_flags::ADDRESS_SPACE_MAX);
+  type_instance_flags new_flags = type->instance_flags ();
+  new_flags.address_space = aspace;
+
+  return make_qualified_type (type, new_flags, nullptr);
+}
+
 /* See gdbtypes.h.  */
 
 type *
@@ -687,6 +706,7 @@ replace_type (struct type *ntype, struct type *type)
 	 symbol readers which do construct address-class variants don't
 	 call replace_type().  */
       gdb_assert (chain->address_class () == 0);
+      gdb_assert (chain->address_space () == 0);
 
       chain->set_length (type->length ());
       chain = chain->chain;
@@ -4996,6 +5016,8 @@ recursive_dump_type (struct type *type, int spaces)
     gdb_puts (" TYPE_DATA_SPACE");
   if (type->address_class () != 0)
     gdb_printf (" TYPE_ADDRESS_CLASS(%u)", type->address_class ());
+  if (type->address_space () != 0)
+    gdb_printf (" TYPE_ADDRESS_SPACE(%s)", pulongest (type->address_space ()));
   if (type->is_restrict ())
     gdb_puts (" TYPE_RESTRICT");
   if (type->is_atomic ())
diff --git a/gdb/gdbtypes.h b/gdb/gdbtypes.h
index dd2d24fa8e2..3b980e83ce5 100644
--- a/gdb/gdbtypes.h
+++ b/gdb/gdbtypes.h
@@ -109,6 +109,7 @@ struct type_instance_flags
 	    && is_volatile == other.is_volatile
 	    && harvard_aspace == other.harvard_aspace
 	    && address_class == other.address_class
+	    && address_space == other.address_space
 	    && is_nottext == other.is_nottext
 	    && is_restrict == other.is_restrict
 	    && is_atomic == other.is_atomic);
@@ -131,6 +132,8 @@ struct type_instance_flags
       harvard_aspace = other.harvard_aspace;
     if (address_class == 0)
       address_class = other.address_class;
+    if (address_space == 0)
+      address_space = other.address_space;
 
     is_nottext = is_nottext || other.is_nottext;
     is_restrict = is_restrict || other.is_restrict;
@@ -156,6 +159,16 @@ struct type_instance_flags
      classes.  */
   unsigned int address_class : 2;
 
+  /* Address space field.  Some architectures have distinct memory
+     storage banks for which the pointer value alone is not sufficient
+     to distinguish which storage is being pointed at.  For these
+     cases, the address space value, which typically comes from the
+     DW_AT_address_space attribute in DWARF, determines where to go.
+     The particular meaning of which actual value represents which
+     storage, is target-specific.  */
+  arch_addr_space_id address_space : 4;
+  static constexpr arch_addr_space_id ADDRESS_SPACE_MAX = 15;
+
   /* Not textual.  By default, GDB treats all single byte integers as
      characters (or elements of strings) unless this flag is set.  */
   bool is_nottext : 1;
@@ -1235,6 +1248,12 @@ struct type
     return this->m_instance_flags.address_class;
   }
 
+  /* Return the address space of this type.  */
+  arch_addr_space_id address_space () const
+  {
+    return (arch_addr_space_id) this->m_instance_flags.address_space;
+  }
+
   /* Get the bounds bounds of this type.  The type must be a range type.  */
   range_bounds *bounds () const
   {
@@ -2493,6 +2512,9 @@ extern struct type *make_type_with_harvard_address_space
 extern struct type *make_type_with_address_class
   (struct type *type, unsigned int address_class);
 
+extern type *make_type_with_address_space (type *type,
+					   arch_addr_space_id aspace);
+
 /* Implement direct support for MEMBER_TYPE in GNU C++.
    TO_TYPE is the type of the member.  DOMAIN is the type of the aggregate that
    the member belongs to.  */
diff --git a/gdb/language.c b/gdb/language.c
index 3d36573a958..6b59785f7dc 100644
--- a/gdb/language.c
+++ b/gdb/language.c
@@ -44,6 +44,7 @@
 #include "c-lang.h"
 #include <algorithm>
 #include "gdbarch.h"
+#include "arch-utils.h"
 
 static void set_range_case (void);
 
@@ -628,7 +629,21 @@ language_defn::watch_location_expression (struct type *type,
   /* Generates an expression that assumes a C like syntax is valid.  */
   type = check_typedef (check_typedef (type)->target_type ());
   std::string name = type_to_string (type);
-  return xstrprintf ("* (%s *) %s", name.c_str (), core_addr_to_string (addr));
+
+  const char *aspace_name = "";
+  const char *aspace_operator = "";
+  arch_addr_space_id aspace
+    = gdbarch_address_space_id_from_core_address (type->arch (), addr);
+
+  if (aspace != 0)
+    {
+      aspace_name = gdbarch_address_space_id_to_name (type->arch (), aspace);
+      aspace_operator = "#";
+      addr = gdbarch_segment_address_from_core_address (type->arch (), addr);
+    }
+
+  return xstrprintf ("* (%s *) %s%s%s", name.c_str (),
+		     aspace_name, aspace_operator, core_addr_to_string (addr));
 }
 
 /* See language.h.  */
diff --git a/gdb/printcmd.c b/gdb/printcmd.c
index bd921b2190d..cb88fdde7e1 100644
--- a/gdb/printcmd.c
+++ b/gdb/printcmd.c
@@ -376,8 +376,10 @@ print_scalar_formatted (const gdb_byte *valaddr, struct type *type,
 
   /* If the value is a pointer, and pointers and addresses are not the
      same, then at this point, the value's length (in target bytes) is
-     gdbarch_addr_bit/TARGET_CHAR_BIT, not type->length ().  */
-  if (type->code () == TYPE_CODE_PTR)
+     gdbarch_addr_bit/TARGET_CHAR_BIT, not type->length ().
+     An exception to this is pointers to address spaces, which may have
+     different sizes.  */
+  if (type->code () == TYPE_CODE_PTR && type->address_space () == 0)
     len = gdbarch_addr_bit (gdbarch) / TARGET_CHAR_BIT;
 
   /* If we are printing it as unsigned, truncate it in case it is actually
diff --git a/gdb/testsuite/gdb.rocm/aspace-pointer.cpp b/gdb/testsuite/gdb.rocm/aspace-pointer.cpp
new file mode 100644
index 00000000000..e7303f64bb4
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/aspace-pointer.cpp
@@ -0,0 +1,68 @@
+/* Copyright (C) 2026 Free Software Foundation, Inc.
+   Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Run a HIP program with pointers to particular address spaces.  */
+
+#include <hip/hip_runtime.h>
+#include "gdb_watchdog.h"
+#include "rocm-test-utils.h"
+
+/* Address space values are defined by the ABI.  */
+#define LOCAL __attribute__((address_space(3)))
+#define PRIVATE_LANE __attribute__((address_space(5)))
+
+__device__ short global_var = 1234;
+
+__global__ void
+kernel ()
+{
+  __shared__ short local_var1;
+  __shared__ short local_var2;
+  local_var1 = 101;
+  local_var2 = 102;
+
+  short priv_var1 = 41;
+  short priv_var2 = 42;
+
+  short *generic_ptr0 = nullptr;
+  short *generic_ptr1 = &global_var;
+  short *generic_ptr2 = &local_var1;
+  short *generic_ptr3 = &priv_var1;
+
+  LOCAL short *local_ptr0 = nullptr;
+  LOCAL short *local_ptr1 = (LOCAL short *) &local_var2;
+  LOCAL short *local_ptr2 = (LOCAL short *) generic_ptr2;
+
+  PRIVATE_LANE short *priv_ptr0 = nullptr;
+  PRIVATE_LANE short *priv_ptr1 = (PRIVATE_LANE short *) &priv_var2;
+  PRIVATE_LANE short *priv_ptr2 = (PRIVATE_LANE short *) generic_ptr3;
+
+  int sizeof_generic_ptr = sizeof(generic_ptr0);
+  int sizeof_local_ptr = sizeof(local_ptr0);
+  int sizeof_priv_ptr = sizeof(priv_ptr0);
+
+  NOP (1); /* break-here */
+}
+
+int
+main ()
+{
+  kernel<<<5, 200>>> ();
+  CHECK (hipDeviceSynchronize ());
+  return 0;
+}
diff --git a/gdb/testsuite/gdb.rocm/aspace-pointer.exp b/gdb/testsuite/gdb.rocm/aspace-pointer.exp
new file mode 100644
index 00000000000..63b5128000a
--- /dev/null
+++ b/gdb/testsuite/gdb.rocm/aspace-pointer.exp
@@ -0,0 +1,91 @@
+# Copyright 2026 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+# Test reading/writing pointers to address spaces.
+
+load_lib rocm.exp
+
+require allow_hip_tests
+
+standard_testfile .cpp
+
+if {[prepare_for_testing "failed to prepare" $testfile $srcfile {debug hip}]} {
+    return
+}
+
+with_rocm_gpu_lock {
+    if {![runto_main]} {
+	return
+    }
+
+    set bp_line [gdb_get_line_number "break-here"]
+    gdb_breakpoint $bp_line -allow-pending -temporary
+    gdb_continue_to_breakpoint "bp"
+
+    with_test_prefix "dereference" {
+	gdb_test "print *generic_ptr1" "= 1234"
+	gdb_test "print *generic_ptr2" "= 101"
+	gdb_test "print *generic_ptr3" "= 41"
+	gdb_test "print *local_ptr1" "= 102"
+	gdb_test "print *local_ptr2" "= 101"
+	gdb_test "print *priv_ptr1" "= 42"
+	gdb_test "print *priv_ptr2" "= 41"
+
+	gdb_test "print *generic_ptr0" "Cannot access memory.*"
+	gdb_test "print *local_ptr0" "Cannot access memory.*"
+	gdb_test "print *priv_ptr0" "Cannot access memory.*"
+    }
+
+    with_test_prefix "sizeof" {
+	gdb_test "print sizeof_generic_ptr == sizeof(generic_ptr0)" "= true"
+	gdb_test "print sizeof_local_ptr == sizeof(local_ptr0)" "= true"
+	gdb_test "print sizeof_priv_ptr == sizeof(priv_ptr0)" "= true"
+    }
+
+    with_test_prefix "illegal assignment" {
+	gdb_test "print local_ptr0 = &global_var" \
+	    "Cannot convert pointer-to-'global' to a pointer-to-'local'"
+	gdb_test "print local_ptr0 = &priv_var1" \
+	    "Cannot convert pointer-to-'private_lane' to a pointer-to-'local'"
+	gdb_test "print local_ptr0 = generic_ptr1" \
+	    "Cannot convert pointer-to-'generic' to a pointer-to-'local'"
+	gdb_test "print priv_ptr0 = &global_var" \
+	    "Cannot convert pointer-to-'global' to a pointer-to-'private_lane'"
+	gdb_test "print priv_ptr0 = &local_var1" \
+	    "Cannot convert pointer-to-'local' to a pointer-to-'private_lane'"
+	gdb_test "print priv_ptr0 = generic_ptr1" \
+	    "Cannot convert pointer-to-'generic' to a pointer-to-'private_lane'"
+    }
+
+    with_test_prefix "legal assignment" {
+	gdb_test "print generic_ptr0 = &local_var1" "generic#0x.*"
+	gdb_test "print *generic_ptr0" "= 101" "after generic <- local"
+	gdb_test "print generic_ptr0 = &priv_var1" "generic#0x.*"
+	gdb_test "print *generic_ptr0" "= 41" "after generic <- priv"
+
+	gdb_test "print local_ptr0 = &local_var2" "local#0x.*"
+	gdb_test "print *local_ptr0" "= 102"
+	gdb_test "print local_ptr0 == local_ptr1" "= true"
+
+	gdb_test "print priv_ptr0 = &priv_var2" "private_lane#0x.*"
+	gdb_test "print *priv_ptr0" "= 42"
+	gdb_test "print priv_ptr0 == priv_ptr1" "= true"
+
+	gdb_test "print generic_ptr0 = local_ptr0" "generic#0x.*"
+	gdb_test "print *generic_ptr0 == *local_ptr0" "= true"
+	gdb_test "print generic_ptr0 = priv_ptr0" "generic#0x.*"
+	gdb_test "print *generic_ptr0 == *priv_ptr0" "= true"
+    }
+}
diff --git a/gdb/testsuite/gdb.rocm/aspace-user-input.exp b/gdb/testsuite/gdb.rocm/aspace-user-input.exp
index df32c4b4a9a..06e40448e9d 100644
--- a/gdb/testsuite/gdb.rocm/aspace-user-input.exp
+++ b/gdb/testsuite/gdb.rocm/aspace-user-input.exp
@@ -102,20 +102,9 @@ gdb_test_no_output "set var \$private_wave_ptr = \(unsigned int \*\) private_wav
 gdb_test "print \$private_wave_ptr" \
 	 " = \\(unsigned int \\*\\) private_wave#0x0"
 
-# Conversion to integral types.
-gdb_test_no_output "set var \$long_long = \(unsigned long long\) private_lane#0x0"
-
-# Two representation of the same test.
-gdb_test_no_output "set var \$generic_ptr = \(unsigned int \*\) \$long_long"
+# Write and read back.
 gdb_test_no_output "set var \*\$private_lane_ptr = 1"
 gdb_test "print \*\$private_lane_ptr" " = 1" "expect one in private_lane"
-gdb_test "print \*\$generic_ptr" " = 1" "expect one in generic"
-
-gdb_test_no_output "set var \*\$generic_ptr = 2"
-gdb_test "print \*\$generic_ptr" " = 2" "expect two in generic"
-gdb_test "print \*\$private_lane_ptr" " = 2" "expect two in private_lane"
-
-gdb_test "print \*\$generic_ptr == \*\$private_lane_ptr" "true"
 
 # Examine memory.
 gdb_test "x/1wx private_lane#0x0" \
@@ -123,6 +112,15 @@ gdb_test "x/1wx private_lane#0x0" \
 gdb_test "with language fortran -- x/1wx private_lane#0x0" \
     "private_lane#0x0:\[ \t\]+$hex"
 
+# Address space propagates with casting.
+gdb_test_multiple "print (int *) local#0x123" "casting" {
+    -re -wrap "\\\$($decimal) = .*" {
+	set last_val $expect_out(1,string)
+	gdb_test "maint print type \$$last_val" \
+	    "TYPE_ADDRESS_SPACE.*"
+    }
+}
+
 if { [istarget "x86_64-*-*"] } {
     # On x86_64, high address (with the MSB bit set - bit 47 for 4 level page
     # tables and bit 56 for 5 level page table) are sign extended.
diff --git a/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp b/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp
index 72b9555886b..9d078b510a4 100644
--- a/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp
+++ b/gdb/testsuite/gdb.rocm/deref-scoped-pointer.exp
@@ -47,14 +47,10 @@ if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
 # the program; we only need to test more than one.  For each lane, the
 # value printed should be the same as the lane number.  If KFAIL is
 # true, expect a failure.
-proc test_print_each_lane {print_what {kfail 0}} {
+proc test_print_each_lane {print_what} {
     for {set i 0} {$i < 2} {incr i} {
 	with_test_prefix "lane $i" {
 	    gdb_test "lane $i" "Switching to .* lane $i .*" "switch lane"
-
-	    if {$kfail} {
-		setup_kfail "*-*-*" "xxxx"
-	    }
 	    gdb_test "print $print_what" " = $i"
 	}
     }
@@ -75,33 +71,15 @@ proc test {} {
     # private_lane address space.  Iterating over lanes and
     # dereferencing that pointer should print a different value per
     # lane, just like the program itself does with printf.
+    #
+    # 'global_ptr2' is a generic null pointer.  Printing it should be
+    # fine.
     with_test_prefix "global pointer set by program" {
-	# This currently yields:
-	#
-	#  (gdb) p global_ptr
-	#  $1 = (size_t *) 0x2000000000008
-	#
-	# But should instead yield:
-	#
-	#  (gdb) p global_ptr
-	#  $1 = (size_t *) generic#0x2000000000008
-	#
-	gdb_test_multiple "print global_ptr" "" {
-	    -re -wrap " = \\\(size_t \\\*\\\) generic#$::hex" {
-		pass $gdb_test_name
-	    }
-	    -re -wrap " = \\\(size_t \\\*\\\) $::hex" {
-		# GDB is printing the generic address bytes raw?
-		setup_kfail "*-*-*" "xxxx"
-		fail $gdb_test_name
-	    }
-	}
+	gdb_test "print global_ptr" \
+	    " = \\\(size_t \\\*\\\) generic#$::hex"
+	test_print_each_lane "*global_ptr"
 
-	# This is kfailed because it currently fails with:
-	#  print *global_ptr
-	#  Cannot access memory at address 0x2000000000008
-	#  (gdb)
-	test_print_each_lane "*global_ptr" 1
+	gdb_test "print global_ptr2" "generic#0x0"
     }
 
     # Test the same, but write to a global pointer in GDB instead of
@@ -110,20 +88,8 @@ proc test {} {
 	# GDB here should understand that 'global_ptr2' is a pointer
 	# to a generic address, and convert appropriately if
 	# necessary.
-	gdb_test_multiple "print global_ptr2 = &local_var" "" {
-	    -re -wrap " = \\\(size_t \\\*\\\) private_lane#$::hex" {
-		# GDB is printing 'global_ptr2' as if it were a
-		# private_lane pointer, which indicates it has failed
-		# to convert to generic address, and likely wrote its
-		# internal representation of a private_lane pointer
-		# over the generic address pointer bytes.
-		setup_kfail "*-*-*" "xxxx"
-		fail $gdb_test_name
-	    }
-	    -re -wrap " = \\\(size_t \\\*\\\) generic#$::hex" {
-		pass $gdb_test_name
-	    }
-	}
+	gdb_test "print global_ptr2 = &local_var" \
+	    " = \\\(size_t \\\*\\\) generic#$::hex"
 	test_print_each_lane "*global_ptr2"
     }
 
@@ -136,10 +102,32 @@ proc test {} {
 	test_print_each_lane "*\$conv_var"
     }
 
-    # Both global_ptr and global_ptr2 should point to the same generic
-    # address, and thus the assertion in the program should not fail.
-    # However, currently it does fail.
-    setup_kfail "*-*-*" "xxxx"
+    # The contents of global pointers should be exactly the same.
+    # Compare the contents,then continue the program.  We should not
+    # hit the assertion.
+    #
+    # Do not use "print" to fetch the contents here, because address
+    # space may be encoded in pointer value, which is hidden by
+    # "print".  We want to check the raw contents.
+    set global_ptr_content "invalid"
+    gdb_test_multiple "x/1gx &global_ptr" "global_ptr contents" {
+	-re -wrap "$::hex <global_ptr>:\[ \t\]+($::hex)" {
+	    set global_ptr_content $expect_out(1,string)
+	    pass $gdb_test_name
+	}
+    }
+
+    set global_ptr2_content "unset"
+    gdb_test_multiple "x/1gx &global_ptr2" "global_ptr2 contents" {
+	-re -wrap "$::hex <global_ptr2>:\[ \t\]+($::hex)" {
+	    set global_ptr2_content $expect_out(1,string)
+	    pass $gdb_test_name
+	}
+    }
+
+    verbose -log "global_ptr=$global_ptr_content, global_ptr2=$global_ptr2_content"
+    gdb_assert {$global_ptr_content == $global_ptr2_content}
+
     gdb_continue_to_end "continue to end" "continue" 1
 }
 
diff --git a/gdb/testsuite/gdb.rocm/finish.exp b/gdb/testsuite/gdb.rocm/finish.exp
index 2f9ce458857..b5b0b6cbde4 100644
--- a/gdb/testsuite/gdb.rocm/finish.exp
+++ b/gdb/testsuite/gdb.rocm/finish.exp
@@ -118,19 +118,19 @@ proc_with_prefix pointer {} {
     gdb_breakpoint "returnPtr if \$_lane == 4"
     gdb_continue_to_breakpoint "returnPtr"
     gdb_test "finish" \
-	"Value returned is \\\$$::decimal = \\(const Custom<char, 8> \\*\\) $::hex <returnPtr\\(\\)::vv>"
+	"Value returned is \\\$$::decimal = \\(const Custom<char, 8> \\*\\) generic#$::hex"
 
     gdb_breakpoint "returnPtr2 if \$_lane == 3"
     with_test_prefix "private address" {
 	gdb_continue_to_breakpoint "returnPtr2"
 	gdb_test "finish" \
-	    "Value returned is \\\$$::decimal = \\(int \\*\\) $::hex"
+	    "Value returned is \\\$$::decimal = \\(int \\*\\) generic#$::hex"
     }
 
     with_test_prefix "global address" {
 	gdb_continue_to_breakpoint "returnPtr2"
 	gdb_test "finish" \
-	    "Value returned is \\\$$::decimal = \\(int \\*\\) $::hex <someGlobal>"
+	    "Value returned is \\\$$::decimal = \\(int \\*\\) generic#$::hex"
     }
 
 }
@@ -139,20 +139,19 @@ proc_with_prefix reference {} {
     gdb_breakpoint "returnRef if \$_lane == 4"
     gdb_continue_to_breakpoint "returnRef"
     gdb_test "finish" \
-	"Value returned is \\\$$::decimal = \\(const Custom<char, 8> &\\) @$::hex: \\{data = \"abcdefgh\"\\}"
+	"Value returned is \\\$$::decimal = \\(const Custom<char, 8> &\\) @generic#$::hex: \\{data = \"abcdefgh\"\\}"
 
     gdb_breakpoint "returnRef2 if \$_lane == 4"
     with_test_prefix "private address" {
 	gdb_continue_to_breakpoint "returnRef2"
-	setup_xfail "*-*-*"
 	gdb_test "finish" \
-	    "Value returned is \\\$$::decimal = \\(int &\\) @$::hex: 4"
+	    "Value returned is \\\$$::decimal = \\(int &\\) @generic#$::hex: 4"
     }
 
     with_test_prefix "global_address" {
 	gdb_continue_to_breakpoint "returnRef2"
 	gdb_test "finish" \
-	    "Value returned is \\\$$::decimal = \\(int &\\) @$::hex: 42"
+	    "Value returned is \\\$$::decimal = \\(int &\\) @generic#$::hex: 42"
     }
 }
 
diff --git a/gdb/testsuite/gdb.rocm/generic-address.cpp b/gdb/testsuite/gdb.rocm/generic-address.cpp
index d0dc708aef4..d26ad6f6b0a 100644
--- a/gdb/testsuite/gdb.rocm/generic-address.cpp
+++ b/gdb/testsuite/gdb.rocm/generic-address.cpp
@@ -43,19 +43,27 @@ kern ()
       assert (local == 1337);
       /* Break here 2: If hit, the assert must have gone smoothly.  */
       NOP (2);
+
+      assert (local == 1991);
+      /* Break here 3: If hit, the assert must have gone smoothly.  */
+      NOP (3);
     }
 
-  /* Break here 3; debugger's actions:
+  /* Break here 4; debugger's actions:
      - Check if still in lane 11.
      - Run some other checks.
      - Change "arr[1..4]" to 0x87654321.  */
-  NOP (3);
+  NOP (4);
 
   if (lane == 11)
     {
       assert (*((uint32_t *) &arr[1]) == 0x87654321);
-      /* Break here 4: If hit, the assert must have gone smoothly.  */
-      NOP (4);
+      /* Break here 5: If hit, the assert must have gone smoothly.  */
+      NOP (5);
+
+      assert (*((uint32_t *) &arr[1]) == 0xfedcba98);
+      /* Break here 6: If hit, the assert must have gone smoothly.  */
+      NOP (6);
     }
 }
 
diff --git a/gdb/testsuite/gdb.rocm/generic-address.exp b/gdb/testsuite/gdb.rocm/generic-address.exp
index 4fdfb800784..0fdb08cacb5 100644
--- a/gdb/testsuite/gdb.rocm/generic-address.exp
+++ b/gdb/testsuite/gdb.rocm/generic-address.exp
@@ -44,7 +44,7 @@ proc do_test {} {
 	}
 
 	# Set breakpoints in device code.  One hit is enough.
-	for {set i 1} {$i <= 4} {incr i} {
+	for {set i 1} {$i <= 6} {incr i} {
 	    set bp_line [gdb_get_line_number "Break here $i"]
 	    gdb_breakpoint $bp_line -allow-pending -temporary
 	}
@@ -68,29 +68,11 @@ proc do_test {} {
 		}
 	    }
 
-	    # N.B.:
-	    # At the time of writing this test, the compiler does not attach
-	    # the DWARF information (DW_AT_LLVM_address_space) to pointer or
-	    # reference DIEs.  Even if it would, rocgdb does not know how to
-	    # use it.  This imposes limitations as it cannot be determined
-	    # where the value of "gen_local_addr" points to or how to handle
-	    # a reference to "local" variable.
-
-	    # For future/backward compatibility, we try to accommodate both
-	    # cases where "gen_local_addr" is a basic global pointer or a
-	    # generic one.
 	    set addr ""
 	    gdb_test_multiple "p gen_local_addr" "extract address type" {
-		-re " = \\(volatile int \\*\\) (generic\\#$::hex)\r\n" {
+		-re -wrap " = \\(volatile int \\*\\) (generic#$::hex)" {
 		    set addr $expect_out(1,string)
-		    exp_continue
-		}
-		-re " = \\(volatile int \\*\\) ($::hex)\r\n" {
-		    set addr "generic#$expect_out(1,string)"
-		    exp_continue
-		}
-		-re "$::gdb_prompt $" {
-		    gdb_assert { $addr != "" } $gdb_test_name
+		    pass $gdb_test_name
 		}
 	    }
 
@@ -106,9 +88,15 @@ proc do_test {} {
 	    gdb_test "print local" " = 1337" "read what was written"
 	    gdb_test "continue" "NOP \\(2\\);" \
 		"modification is reflected in the program"
+
+	    # Write again, this time through the pointer.
+	    gdb_test_no_output "set *gen_local_addr = 1991" "change again"
+	    gdb_test "print local" " = 1991" "read again"
+	    gdb_test "continue" "NOP \\(3\\);" \
+		"modification via pointer is reflected in the program"
 	}
 
-	gdb_test "continue" "NOP \\(3\\);" "look into \"arr\" array"
+	gdb_test "continue" "NOP \\(4\\);" "look into \"arr\" array"
 	gdb_test "info lane" "\\*\[ \t\]+$lane\[ \t\]+A\[ \t\]+.*" \
 	    "must be in the right lane ($lane)"
 
@@ -125,16 +113,9 @@ proc do_test {} {
 
 	    set addr ""
 	    gdb_test_multiple "p gen_arr1_addr" "extract address type" {
-		-re " = \\(volatile uint32_t \\*\\) (generic\\#$::hex)\r\n" {
+		-re -wrap " = \\(volatile uint32_t \\*\\) (generic#$::hex)" {
 		    set addr $expect_out(1,string)
-		    exp_continue
-		}
-		-re " = \\(volatile uint32_t \\*\\) ($::hex)\r\n" {
-		    set addr "generic#$expect_out(1,string)"
-		    exp_continue
-		}
-		-re "$::gdb_prompt $" {
-		    gdb_assert { $addr != "" } $gdb_test_name
+		    pass $gdb_test_name
 		}
 	    }
 
@@ -149,8 +130,15 @@ proc do_test {} {
 	    gdb_test_no_output "set {uint32_t} $addr = 0x87654321" "change"
 	    gdb_test "print /x {uint32_t}(arr + 1)" \
 	      " = 0x87654321" "read what was written"
-	    gdb_test "continue" "NOP \\(4\\);" \
+	    gdb_test "continue" "NOP \\(5\\);" \
 		"modification is reflected in the program"
+
+	    # Write again, this time through the pointer.
+	    gdb_test_no_output "set *gen_arr1_addr = 0xfedcba98" "change again"
+	    gdb_test "print /x {uint32_t}(arr + 1)" \
+	      " = 0xfedcba98" "read again"
+	    gdb_test "continue" "NOP \\(6\\);" \
+		"modification via pointer is reflected in the program"
 	}
     }
 }
diff --git a/gdb/valops.c b/gdb/valops.c
index 2aab81fe870..dd894e7b122 100644
--- a/gdb/valops.c
+++ b/gdb/valops.c
@@ -289,6 +289,43 @@ value_cast_pointers (struct type *type, struct value *arg2,
 {
   struct type *type1 = check_typedef (type);
   struct type *type2 = check_typedef (arg2->type ());
+
+  arch_addr_space_id aspace1 = type1->address_space ();
+  arch_addr_space_id aspace2 = type2->address_space ();
+
+  /* If the value's type has address space information, propagate
+     that information to the result.  This allows doing things like
+
+       (int *) local#123
+       (int *) &a_var_in_local_address_space
+
+     which yield 'int *' with "local" as the address space.  */
+  if (aspace2 != 0 && aspace1 == 0)
+    {
+      type = make_type_with_address_space (type, aspace2);
+      type1 = check_typedef (type);
+      aspace1 = type1->address_space ();
+    }
+
+  if (aspace1 != aspace2
+      && gdbarch_pointer_to_pointer_p (type1->arch ()))
+    {
+      /* This happens when there is, say, a pointer PTR to generic
+	 address space and the user does
+
+	 (gdb) p ptr = private_lane#0x1234.
+
+	 Convert from one address space to the other.  */
+      CORE_ADDR from_addr = value_as_address (arg2);
+      CORE_ADDR to_addr = gdbarch_pointer_to_pointer (type1->arch (), type2,
+						      from_addr, type1);
+
+      /* Convert the pointer type by clearing the address space
+	 and then using the new one.  */
+      type2 = make_type_with_address_space (type2, aspace1);
+      arg2 = value_from_pointer (type2, to_addr);
+    }
+
   struct type *t1 = check_typedef (type1->target_type ());
   struct type *t2 = check_typedef (type2->target_type ());
 
@@ -630,11 +667,10 @@ value_cast (struct type *type, struct value *arg2)
     {
       return value::zero (to_type, not_lval);
     }
+  else if (code1 == TYPE_CODE_PTR && code2 == TYPE_CODE_PTR)
+    return value_cast_pointers (to_type, arg2, 0);
   else if (type->length () == type2->length ())
     {
-      if (code1 == TYPE_CODE_PTR && code2 == TYPE_CODE_PTR)
-	return value_cast_pointers (to_type, arg2, 0);
-
       arg2 = arg2->copy ();
 
       struct type *resolved_type
diff --git a/gdb/value.c b/gdb/value.c
index 38eadabba23..e121ca1bef3 100644
--- a/gdb/value.c
+++ b/gdb/value.c
@@ -3657,6 +3657,11 @@ value_from_mpz (struct type *type, const gdb_mpz &v)
 struct value *
 value_from_pointer (struct type *type, CORE_ADDR addr)
 {
+  arch_addr_space_id aspace
+    = gdbarch_address_space_id_from_core_address (type->arch (), addr);
+  if (aspace != 0)
+    type = make_type_with_address_space (type, aspace);
+
   struct value *val = value::allocate (type);
 
   store_typed_address (val->contents_raw ().data (),
diff --git a/include/dwarf2.def b/include/dwarf2.def
index dbff7a5a02d..2bf682ec0ef 100644
--- a/include/dwarf2.def
+++ b/include/dwarf2.def
@@ -488,6 +488,7 @@ DW_AT (DW_AT_PGI_lstride, 0x3a02)
 /* LLVM extensions for heterogeneous targets.  See
    https://llvm.org/docs/AMDGPUDwarfExtensionsForHeterogeneousDebugging.html.  */
 DW_AT (DW_AT_LLVM_lane_pc, 0x3e0b)
+DW_AT (DW_AT_LLVM_address_space, 0x3e10)
 /* Apple extensions.  */
 DW_AT (DW_AT_APPLE_optimized, 0x3fe1)
 DW_AT (DW_AT_APPLE_flags, 0x3fe2)
diff --git a/include/dwarf2.h b/include/dwarf2.h
index 7f43ad0835e..e34fe885345 100644
--- a/include/dwarf2.h
+++ b/include/dwarf2.h
@@ -617,6 +617,22 @@ enum dwarf_sect_v5
   DW_SECT_MAX_V5 = 8
 };
 
+/* Codes for DW_AT_address_space attribute.
+   See https://dwarfstd.org/issues/260211.1.html.  */
+enum dwarf_address_space
+{
+  DW_ASPACE_default = 0,
+
+  /* AMD GPU address spaces.
+     See https://llvm.org/docs/AMDGPUUsage.html#address-space-identifier.  */
+  DW_ASPACE_LLVM_AMDGPU_generic = 1,
+  DW_ASPACE_LLVM_AMDGPU_region = 2,
+  DW_ASPACE_LLVM_AMDGPU_local = 3,
+  /* Reserved.  */
+  DW_ASPACE_LLVM_AMDGPU_private_lane = 5,
+  DW_ASPACE_LLVM_AMDGPU_private_wave = 6,
+};
+
 #ifdef __cplusplus
 extern "C" {
 #endif /* __cplusplus */