Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
gmarkall merged 8 commits into from
Dec 19, 2025
Merged

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

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
c8e4a8d
Remove the dwarfAddressSpace attribute from DIDerivedType for shared …
jiel-nv Dec 16, 2025
9b46894
DCE cleanup since customized address space tracking is no longer needed.
jiel-nv Dec 17, 2025
ab37265
Merge branch 'main' into shared-pointer
jiel-nv Dec 17, 2025
31dde8a
Removed unused import
jiel-nv Dec 17, 2025
7a76f2f
Merge branch 'main' into shared-pointer
jiel-nv Dec 17, 2025
17ecf4b
Merge branch 'main' into shared-pointer
jiel-nv Dec 18, 2025
cd4edef
Merge branch 'main' into shared-pointer
jiel-nv Dec 19, 2025
71da695
Minor tidy-up
gmarkall Dec 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
89 changes: 10 additions & 79 deletions numba_cuda/numba/cuda/debuginfo.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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)

Expand Down Expand Up @@ -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,
Expand Down
25 changes: 0 additions & 25 deletions numba_cuda/numba/cuda/lowering.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -1684,31 +1683,13 @@ 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):
Copy link
Contributor

@gmarkall gmarkall Dec 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this function can now be deleted as it entirely defers to Lower._lower_call_normal().

Copy link
Contributor Author

@jiel-nv jiel-nv Dec 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for catching this!

# 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]
Expand Down Expand Up @@ -1834,12 +1815,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 = {}
Expand Down
19 changes: 12 additions & 7 deletions numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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),)

Expand Down
Loading