Remove customized address space tracking and address class emission in debug info (#669)

Fix issue [#627](#627)

Remove the dwarfAddressSpace attribute from DIDerivedType for shared
memory pointers;
DCE cleanup since customized address space tracking is no longer needed.

The Enum data structure of DwarfAddressClass and the corrsponding
mapping function (map NVVM address space to DWARF address class) are
intentionly retained for potential future usage.

---------

Co-authored-by: Graham Markall <gmarkall@nvidia.com>

Author: jiel-nv

numba_cuda/numba/cuda/debuginfo.py

@@ -646,11 +646,6 @@ def __init__(self, module, filepath, cgctx, directives_only):
         super().__init__(module, filepath, cgctx, directives_only)
         # Cache for local variable metadata type and line deduplication
         self._vartypelinemap = {}
-        # Variable address space dictionary
-        self._var_addrspace_map = {}
-
-    def _set_addrspace_map(self, map):
-        self._var_addrspace_map = map
 
     def _var_type(self, lltype, size, datamodel=None):
         is_bool = False
@@ -826,64 +821,6 @@ def _var_type(self, lltype, size, datamodel=None):
                     is_distinct=True,
                 )
 
-        # Check if there's actually address space info to handle
-        addrspace = getattr(self, "_addrspace", None)
-        if (
-            isinstance(lltype, ir.LiteralStructType)
-            and datamodel is not None
-            and datamodel.inner_models()
-            and addrspace not in (None, 0)
-        ):
-            # Process struct with datamodel that has address space info
-            meta = []
-            offset = 0
-            for element, field, model in zip(
-                lltype.elements, datamodel._fields, datamodel.inner_models()
-            ):
-                size_field = self.cgctx.get_abi_sizeof(element)
-                if isinstance(element, ir.PointerType) and field == "data":
-                    # Create pointer type with correct address space
-                    pointee_size = self.cgctx.get_abi_sizeof(element.pointee)
-                    pointee_model = getattr(model, "_pointee_model", None)
-                    pointee_type = self._var_type(
-                        element.pointee, pointee_size, datamodel=pointee_model
-                    )
-                    meta_ptr = {
-                        "tag": ir.DIToken("DW_TAG_pointer_type"),
-                        "baseType": pointee_type,
-                        "size": _BYTE_SIZE * size_field,
-                    }
-                    dwarf_addrclass = self.get_dwarf_address_class(addrspace)
-                    if dwarf_addrclass is not None:
-                        meta_ptr["dwarfAddressSpace"] = int(dwarf_addrclass)
-                    basetype = m.add_debug_info("DIDerivedType", meta_ptr)
-                else:
-                    basetype = self._var_type(
-                        element, size_field, datamodel=model
-                    )
-                derived_type = m.add_debug_info(
-                    "DIDerivedType",
-                    {
-                        "tag": ir.DIToken("DW_TAG_member"),
-                        "name": field,
-                        "baseType": basetype,
-                        "size": _BYTE_SIZE * size_field,
-                        "offset": offset,
-                    },
-                )
-                meta.append(derived_type)
-                offset += _BYTE_SIZE * size_field
-
-            return m.add_debug_info(
-                "DICompositeType",
-                {
-                    "tag": ir.DIToken("DW_TAG_structure_type"),
-                    "name": f"{datamodel.fe_type}",
-                    "elements": m.add_metadata(meta),
-                    "size": offset,
-                },
-                is_distinct=True,
-            )
         # For other cases, use upstream Numba implementation
         return super()._var_type(lltype, size, datamodel=datamodel)
 
@@ -936,22 +873,16 @@ def mark_variable(
                 # to llvm.dbg.value
                 return
             else:
-                # Look up address space for this variable
-                self._addrspace = self._var_addrspace_map.get(name)
-                try:
-                    return super().mark_variable(
-                        builder,
-                        allocavalue,
-                        name,
-                        lltype,
-                        size,
-                        line,
-                        datamodel,
-                        argidx,
-                    )
-                finally:
-                    # Clean up address space info
-                    self._addrspace = None
+                return super().mark_variable(
+                    builder,
+                    allocavalue,
+                    name,
+                    lltype,
+                    size,
+                    line,
+                    datamodel,
+                    argidx,
+                )
 
     def update_variable(
         self,

numba_cuda/numba/cuda/lowering.py

@@ -11,7 +11,6 @@
 from numba.cuda import HAS_NUMBA
 from numba.cuda.core import ir
 from numba.cuda import debuginfo, cgutils, utils, typing, types
-from numba import cuda
 from numba.cuda.core import (
     ir_utils,
     targetconfig,
@@ -1684,31 +1683,10 @@ def decref(self, typ, val):
 
 
 class CUDALower(Lower):
-    def _is_shared_array_call(self, fnty):
-        # Check if function type is a cuda.shared.array call
-        if not hasattr(fnty, "typing_key"):
-            return False
-        return fnty.typing_key is cuda.shared.array
-
-    def _lower_call_normal(self, fnty, expr, signature):
-        # Set flag for subsequent store to track shared address space
-        if self.context.enable_debuginfo and self._is_shared_array_call(fnty):
-            self._pending_shared_store = True
-
-        return super()._lower_call_normal(fnty, expr, signature)
-
     def storevar(self, value, name, argidx=None):
         """
         Store the value into the given variable.
         """
-        # Track address space for debug info
-        if self.context.enable_debuginfo and self._pending_shared_store:
-            from numba.cuda.cudadrv import nvvm
-
-            self._addrspace_map[name] = nvvm.ADDRSPACE_SHARED
-            if not name.startswith("$") and not name.startswith("."):
-                self._pending_shared_store = False
-
         # Handle polymorphic variables with CUDA_DEBUG_POLY enabled
         if config.CUDA_DEBUG_POLY:
             src_name = name.split(".")[0]
@@ -1834,12 +1812,6 @@ def pre_lower(self):
         """
         super().pre_lower()
 
-        # Track address space for debug info
-        self._addrspace_map = {}
-        self._pending_shared_store = False
-        if self.context.enable_debuginfo:
-            self.debuginfo._set_addrspace_map(self._addrspace_map)
-
         # Track polymorphic variables for debug info
         self.poly_var_typ_map = {}
         self.poly_var_loc_map = {}

numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py

@@ -885,29 +885,34 @@ def foo():
         """,
         )
 
-    # shared_arr -> composite -> elements[4] (data field at index 4) -> pointer with dwarfAddressSpace: 8
-    # local_arr -> composite -> elements[4] (data field at index 4) -> pointer without dwarfAddressSpace: 8
+    # shared_arr -> composite -> elements[4] (data field at index 4) -> pointer without dwarfAddressSpace
+    # local_arr -> composite -> elements[4] (data field at index 4) -> pointer without dwarfAddressSpace
+    # Note: Shared memory pointers don't have dwarfAddressSpace because they are
+    # cast to generic address space via addrspacecast in cudaimpl.py
     address_class_filechecks = r"""
         CHECK-DAG: [[SHARED_VAR:![0-9]+]] = !DILocalVariable({{.*}}name: "shared_arr"{{.*}}type: [[SHARED_COMPOSITE:![0-9]+]]
         CHECK-DAG: [[SHARED_COMPOSITE]] = {{.*}}!DICompositeType(elements: [[SHARED_ELEMENTS:![0-9]+]]
         CHECK-DAG: [[SHARED_ELEMENTS]] = !{{{.*}}, {{.*}}, {{.*}}, {{.*}}, [[SHARED_DATA:![0-9]+]], {{.*}}, {{.*}}}
         CHECK-DAG: [[SHARED_DATA]] = !DIDerivedType(baseType: [[SHARED_PTR:![0-9]+]], name: "data"
-        CHECK-DAG: [[SHARED_PTR]] = !DIDerivedType({{.*}}dwarfAddressSpace: 8{{.*}}tag: DW_TAG_pointer_type
+        CHECK-DAG: [[SHARED_PTR]] = !DIDerivedType({{.*}}tag: DW_TAG_pointer_type
+        CHECK-NOT: [[SHARED_PTR]]{{.*}}dwarfAddressSpace
 
         CHECK-DAG: [[LOCAL_VAR:![0-9]+]] = !DILocalVariable({{.*}}name: "local_arr"{{.*}}type: [[LOCAL_COMPOSITE:![0-9]+]]
         CHECK-DAG: [[LOCAL_COMPOSITE]] = {{.*}}!DICompositeType(elements: [[LOCAL_ELEMENTS:![0-9]+]]
         CHECK-DAG: [[LOCAL_ELEMENTS]] = !{{{.*}}, {{.*}}, {{.*}}, {{.*}}, [[LOCAL_DATA:![0-9]+]], {{.*}}, {{.*}}}
         CHECK-DAG: [[LOCAL_DATA]] = !DIDerivedType(baseType: [[LOCAL_PTR:![0-9]+]], name: "data"
         CHECK-DAG: [[LOCAL_PTR]] = !DIDerivedType(baseType: {{.*}}tag: DW_TAG_pointer_type
-        CHECK-NOT: [[LOCAL_PTR]]{{.*}}dwarfAddressSpace: 8
+        CHECK-NOT: [[LOCAL_PTR]]{{.*}}dwarfAddressSpace
     """
 
     def _test_shared_memory_address_class(self, dtype):
         """Test that shared memory arrays have correct DWARF address class.
 
-        Shared memory pointers should have addressClass: 8 (DW_AT_address_class
-        for CUDA shared memory) in their debug metadata, while regular local
-        arrays should not have this annotation.
+        Shared memory pointers should NOT have dwarfAddressSpace attribute
+        because they are cast to generic address space via addrspacecast.
+        The runtime pointer type is generic, not shared, so cuda-gdb can
+        correctly dereference them. Local arrays also should not have this
+        attribute.
         """
         sig = (numpy_support.from_dtype(dtype),)