Skip to content

Conversation

@mmason-nvidia
Copy link
Contributor

This is a direct copy of the current version of Numba's GDB pretty-printer to use as a starting point for CUDA specific enhancements.

Closes #691

This is a direct copy of the current version of the Numba GDB pretty-printer to use as a starting point for CUDA specific enhancements.
@copy-pr-bot
Copy link

copy-pr-bot bot commented Jan 5, 2026

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@greptile-apps
Copy link
Contributor

greptile-apps bot commented Jan 5, 2026

Greptile Overview

Greptile Summary

This PR adds a GDB pretty-printer extension for Numba types, copied from upstream Numba as a starting point for CUDA-specific enhancements. The implementation provides formatted output for arrays, complex numbers, tuples, and unicode strings during GDB debugging sessions.

Critical Issues Found:

The code contains 11 critical logic errors that will cause crashes during GDB debugging:

  1. Three unchecked regex matches (lines 59, 62, 85) that will throw AttributeError when type strings don't match expected patterns
  2. Multiple array index access violations (lines 69, 78, 94, 145, 157) that will throw IndexError when data structures are empty or malformed
  3. Unsafe memory operations (lines 107, 174) with inadequate validation that could read invalid memory ranges
  4. Incorrect unicode decoding (line 176) that doesn't respect the kind parameter indicating character width (1/2/4 bytes)

Impact:

These errors will cause the GDB pretty-printer to crash when encountering:

  • Type strings with unexpected formats
  • Empty or malformed array/tuple structures
  • Invalid memory addresses or sizes
  • Unicode strings with multi-byte characters

Since this is intended to be a debugging aid, crashes in the pretty-printer will severely impact the debugging experience. All identified issues should be fixed before merging to ensure reliable GDB integration.

Confidence Score: 1/5

  • This PR contains multiple critical bugs that will cause crashes during debugging sessions
  • Score of 1 reflects the presence of 11 logic errors (8 critical, 3 high-priority) that will cause AttributeError, IndexError, or invalid memory access exceptions. These are not potential issues - they are definite bugs that will trigger when the pretty-printer encounters unexpected but valid input formats. The code lacks essential error handling throughout, making it unsuitable for production use without fixes.
  • misc/gdb_print_extension.py requires immediate attention - every printer class has critical bugs

Important Files Changed

File Analysis

Filename Score Overview
misc/gdb_print_extension.py 1/5 New GDB pretty-printer with multiple critical unchecked regex matches and array access errors that will cause crashes

@greptile-apps
Copy link
Contributor

greptile-apps bot commented Jan 5, 2026

Greptile's behavior is changing!

From now on, if a review finishes with no comments, we will not post an additional "statistics" comment to confirm that our review found nothing to comment on. However, you can confirm that we reviewed your changes in the status check section.

This feature can be toggled off in your Code Review Settings by deselecting "Create a status check for each PR".

@gmarkall
Copy link
Contributor

/ok to test

@copy-pr-bot
Copy link

copy-pr-bot bot commented Jan 12, 2026

/ok to test

@gmarkall, there was an error processing your request: E1

See the following link for more information: https://docs.gha-runners.nvidia.com/cpr/e/1/

@gmarkall
Copy link
Contributor

/ok to test ebfe957

Copy link
Contributor

@greptile-apps greptile-apps bot left a comment

Choose a reason for hiding this comment

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

No files reviewed, no comments

Edit Code Review Agent Settings | Greptile

@gmarkall
Copy link
Contributor

/ok to test 47827ea

Copy link
Contributor

@greptile-apps greptile-apps bot left a comment

Choose a reason for hiding this comment

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

8 files reviewed, 8 comments

Edit Code Review Agent Settings | Greptile

ty_str = ty_str.replace("unaligned ", "").strip()
matcher = re.compile(r"array\((Record.*), (.*), (.*)\)\ \(.*")
# NOTE: need to deal with "Alignment" else dtype size is wrong
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
Copy link
Contributor

Choose a reason for hiding this comment

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

The matcher.match(ty_str) can return None if the regex pattern doesn't match the type string, which will cause an AttributeError when calling .groups() on None. This should be checked before accessing groups.

Suggested change
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
match_result = matcher.match(ty_str)
if match_result is None:
raise ValueError(f"Unable to parse array type string: {ty_str}")
arr_info = [x.strip() for x in match_result.groups()]

arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
dtype_str, ndim_str, order_str = arr_info
rstr = r"Record\((.*\[.*\]);([0-9]+);(True|False)"
rstr_match = re.match(rstr, dtype_str)
Copy link
Contributor

Choose a reason for hiding this comment

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

The re.match(rstr, dtype_str) can return None if the pattern doesn't match, causing an AttributeError on the next line when calling .groups(). Add a check before accessing the match result.

Suggested change
rstr_match = re.match(rstr, dtype_str)
rstr_match = re.match(rstr, dtype_str)
if rstr_match is None:
raise ValueError(f"Unable to parse Record dtype string: {dtype_str}")

dtype_str = struct_entries
else: # simple type
matcher = re.compile(r"array\((.*),(.*),(.*)\)\ \(.*")
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
Copy link
Contributor

Choose a reason for hiding this comment

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

Same issue as line 59 - matcher.match(ty_str) can return None, causing an AttributeError when calling .groups(). This needs a null check.

Suggested change
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
match_result = matcher.match(ty_str)
if match_result is None:
raise ValueError(f"Unable to parse array type string: {ty_str}")
arr_info = [x.strip() for x in match_result.groups()]

def dwarr2inttuple(dwarr):
# Converts a gdb handle to a dwarf array to a tuple of ints
fields = dwarr.type.fields()
lo, hi = fields[0].type.range()
Copy link
Contributor

Choose a reason for hiding this comment

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

Accessing fields[0] without checking if the fields list is empty will cause an IndexError if there are no fields. Add a check to ensure fields is not empty.

Suggested change
lo, hi = fields[0].type.range()
fields = dwarr.type.fields()
if not fields:
raise ValueError("DWARF array type has no fields")
lo, hi = fields[0].type.range()

def to_string(self):
# unituples are arrays
fields = self.val.type.fields()
lo, hi = fields[0].type.range()
Copy link
Contributor

Choose a reason for hiding this comment

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

Same issue as line 94 - accessing fields[0] without verifying the fields list is non-empty will cause an IndexError. Add validation before accessing.

Suggested change
lo, hi = fields[0].type.range()
fields = self.val.type.fields()
if not fields:
raise ValueError("UniTuple type has no fields")
lo, hi = fields[0].type.range()

self.val = val

def to_string(self):
return "%s+%sj" % (self.val["real"], self.val["imag"])
Copy link
Contributor

Choose a reason for hiding this comment

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

The complex number formatting doesn't handle negative imaginary parts correctly. A complex number like 3-5j will be formatted as 3+-5j. Consider using conditional logic to format negative imaginary parts properly.

Suggested change
return "%s+%sj" % (self.val["real"], self.val["imag"])
imag_val = self.val["imag"]
if float(str(imag_val)) < 0:
return "%s%sj" % (self.val["real"], imag_val)
else:
return "%s+%sj" % (self.val["real"], imag_val)

if data != NULL:
# This needs sorting out, encoding is wrong
this_proc = gdb.selected_inferior()
mem = this_proc.read_memory(int(data), nitems * kind)
Copy link
Contributor

Choose a reason for hiding this comment

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

The memory read size calculation nitems * kind has no bounds checking or validation. If nitems or kind contain corrupted or extremely large values from memory, this could attempt to read gigabytes of memory, causing GDB to hang or crash. Consider adding a reasonable upper bound check.

Suggested change
mem = this_proc.read_memory(int(data), nitems * kind)
mem_size = nitems * kind
MAX_STRING_SIZE = 1024 * 1024 # 1MB reasonable limit
if mem_size > MAX_STRING_SIZE:
return "'[String too large: %d bytes]'" % mem_size
mem = this_proc.read_memory(int(data), mem_size)

Comment on lines +179 to +182
if isinstance(mem, memoryview):
buf = bytes(mem).decode()
else:
buf = mem.decode("utf-8")
Copy link
Contributor

Choose a reason for hiding this comment

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

The encoding handling is inconsistent between the two branches. When mem is a memoryview, bytes(mem).decode() is called without specifying an encoding (defaults to UTF-8), but in the else branch, mem.decode("utf-8") explicitly specifies UTF-8. The comment on line 175 also acknowledges "encoding is wrong". Consider standardizing this to explicitly use the correct encoding based on the kind field (which should indicate the character width: 1, 2, or 4 bytes for UTF-8, UTF-16, or UTF-32).

Suggested change
if isinstance(mem, memoryview):
buf = bytes(mem).decode()
else:
buf = mem.decode("utf-8")
# Decode based on kind: 1=UTF-8, 2=UTF-16, 4=UTF-32
encoding_map = {1: 'utf-8', 2: 'utf-16', 4: 'utf-32'}
encoding = encoding_map.get(int(kind), 'utf-8')
if isinstance(mem, memoryview):
buf = bytes(mem).decode(encoding)
else:
buf = mem.decode(encoding)

@gmarkall
Copy link
Contributor

/ok to test df30db0

Copy link
Contributor

@greptile-apps greptile-apps bot left a comment

Choose a reason for hiding this comment

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

12 files reviewed, 12 comments

Edit Code Review Agent Settings | Greptile

ty_str = ty_str.replace("unaligned ", "").strip()
matcher = re.compile(r"array\((Record.*), (.*), (.*)\)\ \(.*")
# NOTE: need to deal with "Alignment" else dtype size is wrong
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
Copy link
Contributor

Choose a reason for hiding this comment

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

the regex match result is not checked for None before calling .groups(), which will cause an AttributeError if the type string doesn't match the expected pattern

Suggested change
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
match = matcher.match(ty_str)
if match is None:
raise ValueError(f"Type string does not match expected pattern: {ty_str}")
arr_info = [x.strip() for x in match.groups()]

Comment on lines +62 to +64
rstr_match = re.match(rstr, dtype_str)
# balign is unused, it's the alignment
fields, balign, is_aligned_str = rstr_match.groups()
Copy link
Contributor

Choose a reason for hiding this comment

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

the regex match result is not checked for None before calling .groups(), which will cause an AttributeError if dtype_str doesn't match the expected Record pattern

Suggested change
rstr_match = re.match(rstr, dtype_str)
# balign is unused, it's the alignment
fields, balign, is_aligned_str = rstr_match.groups()
rstr_match = re.match(rstr, dtype_str)
if rstr_match is None:
raise ValueError(f"dtype_str does not match expected Record pattern: {dtype_str}")
# balign is unused, it's the alignment
fields, balign, is_aligned_str = rstr_match.groups()

dtype_str = struct_entries
else: # simple type
matcher = re.compile(r"array\((.*),(.*),(.*)\)\ \(.*")
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
Copy link
Contributor

Choose a reason for hiding this comment

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

the regex match result is not checked for None before calling .groups(), which will cause an AttributeError if the type string doesn't match the expected pattern

Suggested change
arr_info = [x.strip() for x in matcher.match(ty_str).groups()]
match = matcher.match(ty_str)
if match is None:
raise ValueError(f"Type string does not match expected pattern: {ty_str}")
arr_info = [x.strip() for x in match.groups()]

dt_part = dt_part[0]
if "nestedarray" in dt_part:
raise TypeError("Unsupported sub-type: %s" % f)
dt_as_str = dt_part.split(";")[0].split("=")[1]
Copy link
Contributor

Choose a reason for hiding this comment

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

array index access without bounds checking - if the split operations don't produce the expected number of elements, this will cause an IndexError

Suggested change
dt_as_str = dt_part.split(";")[0].split("=")[1]
parts = dt_part.split(";")
if len(parts) < 1:
raise ValueError(f"Invalid dt_part format: {dt_part}")
type_parts = parts[0].split("=")
if len(type_parts) < 2:
raise ValueError(f"Invalid type format: {parts[0]}")
dt_as_str = type_parts[1]

def dwarr2inttuple(dwarr):
# Converts a gdb handle to a dwarf array to a tuple of ints
fields = dwarr.type.fields()
lo, hi = fields[0].type.range()
Copy link
Contributor

Choose a reason for hiding this comment

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

accessing fields[0] without checking if the fields list is empty will cause an IndexError

Suggested change
lo, hi = fields[0].type.range()
fields = dwarr.type.fields()
if not fields:
raise ValueError("DWARF array type has no fields")
lo, hi = fields[0].type.range()

this_proc = gdb.selected_inferior()
mem = this_proc.read_memory(int(data), nitems * kind)
if isinstance(mem, memoryview):
buf = bytes(mem).decode()
Copy link
Contributor

Choose a reason for hiding this comment

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

calling .decode() without specifying an encoding will default to 'utf-8', which is correct, but it's better to be explicit. Additionally, the kind parameter (which represents the character width: 1, 2, or 4 bytes) is not being used properly to determine the encoding

Suggested change
buf = bytes(mem).decode()
if kind == 1:
buf = bytes(mem).decode('latin-1')
elif kind == 2:
buf = bytes(mem).decode('utf-16')
elif kind == 4:
buf = bytes(mem).decode('utf-32')
else:
buf = bytes(mem).decode('utf-8')

Comment on lines +107 to +110
shp_arr = np.array([max(0, x - 1) for x in shape])
strd_arr = np.array(strides)
extent = np.sum(shp_arr * strd_arr)
extent += int(itemsize)
Copy link
Contributor

Choose a reason for hiding this comment

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

the extent calculation could produce a negative value or zero if shape contains values <= 0, which would cause issues with read_memory. The max(0, x - 1) protects against negative intermediate values but doesn't validate the final extent

Suggested change
shp_arr = np.array([max(0, x - 1) for x in shape])
strd_arr = np.array(strides)
extent = np.sum(shp_arr * strd_arr)
extent += int(itemsize)
# The data extent in bytes is:
# sum(shape * strides)
# get the data, then wire to as_strided
shp_arr = np.array([max(0, x - 1) for x in shape])
strd_arr = np.array(strides)
extent = np.sum(shp_arr * strd_arr)
extent += int(itemsize)
if extent <= 0:
raise ValueError(f"Invalid extent calculated: {extent}")

Comment on lines +69 to +71
splitted = f.split("[")
name = splitted[0]
dt_part = splitted[1:]
Copy link
Contributor

Choose a reason for hiding this comment

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

if the field string f doesn't contain a '[' character, the split will return a single element, but the code accesses splitted[1:] which could be empty

Suggested change
splitted = f.split("[")
name = splitted[0]
dt_part = splitted[1:]
splitted = f.split("[")
if len(splitted) < 2:
raise ValueError(f"Invalid field format, missing '[': {f}")
name = splitted[0]
dt_part = splitted[1:]

return "array([...], dtype=%s, shape=%s)" % (dtype_str, shape)
else:
# Not yet initialized or NULLed out data
buf = list(["NULL/Uninitialized"])
Copy link
Contributor

Choose a reason for hiding this comment

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

use a list literal directly instead of wrapping it in list()

Suggested change
buf = list(["NULL/Uninitialized"])
buf = ["NULL/Uninitialized"]

Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!

self.val = val

def to_string(self):
fields = self.val.type.fields()
Copy link
Contributor

Choose a reason for hiding this comment

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

accessing self.val.type.fields() without checking if it returns an empty list could cause issues in the list comprehension on line 146 if there are no fields

Suggested change
fields = self.val.type.fields()
fields = self.val.type.fields()
if not fields:
return "()"
buf = [str(self.val[f.name]) for f in fields]

@gmarkall gmarkall merged commit 622b678 into NVIDIA:main Jan 14, 2026
89 checks passed
gmarkall added a commit to gmarkall/numba-cuda that referenced this pull request Jan 27, 2026
- Add Python 3.14 to the wheel publishing matrix (NVIDIA#750)
- feat: swap out internal device array usage with `StridedMemoryView` (NVIDIA#703)
- Fix max block size computation in `forall` (NVIDIA#744)
- Fix prologue debug line info pointing to decorator instead of def line (NVIDIA#746)
- Fix kernel return type in DISubroutineType debug metadata (NVIDIA#745)
- Fix missing line info in Jupyter notebooks (NVIDIA#742)
- Fix: Pass correct flags to linker when debugging in the presence of LTOIR code (NVIDIA#698)
- chore(deps): add cuda-pathfinder to pixi deps (NVIDIA#741)
- fix: enable flake8-bugbear lints and fix found problems (NVIDIA#708)
- fix: Fix race condition in CUDA Simulator (NVIDIA#690)
- ci: run tests in parallel (NVIDIA#740)
- feat: users can pass `shared_memory_carveout` to @cuda.jit (NVIDIA#642)
- Fix compatibility with NumPy 2.4: np.trapz and np.in1d removed (NVIDIA#739)
- Pass the -numba-debug flag to libnvvm (NVIDIA#681)
- ci: remove rapids containers from conda ci (NVIDIA#737)
- Use `pathfinder` for dynamic libraries (NVIDIA#308)
- CI: Add CUDA 13.1 testing support (NVIDIA#705)
- Adding `pixi run test` and `pixi run test-par` support (NVIDIA#724)
- Disable per-PR nvmath tests + follow same test practice (NVIDIA#723)
- chore(deps): regenerate pixi lockfile (NVIDIA#722)
- Fix DISubprogram line number to point to function definition line (NVIDIA#695)
- revert: chore(dev): build pixi using rattler (NVIDIA#713) (NVIDIA#719)
- [feat] Initial version of the Numba CUDA GDB pretty-printer (NVIDIA#692)
- chore(dev): build pixi using rattler (NVIDIA#713)
- build(deps): bump the actions-monthly group across 1 directory with 8 updates (NVIDIA#704)
@gmarkall gmarkall mentioned this pull request Jan 27, 2026
kkraus14 pushed a commit that referenced this pull request Jan 28, 2026
- Add Python 3.14 to the wheel publishing matrix (#750)
- feat: swap out internal device array usage with `StridedMemoryView`
(#703)
- Fix max block size computation in `forall` (#744)
- Fix prologue debug line info pointing to decorator instead of def line
(#746)
- Fix kernel return type in DISubroutineType debug metadata (#745)
- Fix missing line info in Jupyter notebooks (#742)
- Fix: Pass correct flags to linker when debugging in the presence of
LTOIR code (#698)
- chore(deps): add cuda-pathfinder to pixi deps (#741)
- fix: enable flake8-bugbear lints and fix found problems (#708)
- fix: Fix race condition in CUDA Simulator (#690)
- ci: run tests in parallel (#740)
- feat: users can pass `shared_memory_carveout` to @cuda.jit (#642)
- Fix compatibility with NumPy 2.4: np.trapz and np.in1d removed (#739)
- Pass the -numba-debug flag to libnvvm (#681)
- ci: remove rapids containers from conda ci (#737)
- Use `pathfinder` for dynamic libraries (#308)
- CI: Add CUDA 13.1 testing support (#705)
- Adding `pixi run test` and `pixi run test-par` support (#724)
- Disable per-PR nvmath tests + follow same test practice (#723)
- chore(deps): regenerate pixi lockfile (#722)
- Fix DISubprogram line number to point to function definition line
(#695)
- revert: chore(dev): build pixi using rattler (#713) (#719)
- [feat] Initial version of the Numba CUDA GDB pretty-printer (#692)
- chore(dev): build pixi using rattler (#713)
- build(deps): bump the actions-monthly group across 1 directory with 8
updates (#704)

<!--

Thank you for contributing to numba-cuda :)

Here are some guidelines to help the review process go smoothly.

1. Please write a description in this text box of the changes that are
being
   made.

2. Please ensure that you have written units tests for the changes
made/features
   added.

3. If you are closing an issue please use one of the automatic closing
words as
noted here:
https://help.github.com/articles/closing-issues-using-keywords/

4. If your pull request is not ready for review but you want to make use
of the
continuous integration testing facilities please label it with `[WIP]`.

5. If your pull request is ready to be reviewed without requiring
additional
work on top of it, then remove the `[WIP]` label (if present) and
replace
it with `[REVIEW]`. If assistance is required to complete the
functionality,
for example when the C/C++ code of a feature is complete but Python
bindings
are still required, then add the label `[HELP-REQ]` so that others can
triage
and assist. The additional changes then can be implemented on top of the
same PR. If the assistance is done by members of the rapidsAI team, then
no
additional actions are required by the creator of the original PR for
this,
otherwise the original author of the PR needs to give permission to the
person(s) assisting to commit to their personal fork of the project. If
that
doesn't happen then a new PR based on the code of the original PR can be
opened by the person assisting, which then will be the PR that will be
   merged.

6. Once all work has been done and review has taken place please do not
add
features or make changes out of the scope of those requested by the
reviewer
(doing this just add delays as already reviewed code ends up having to
be
re-reviewed/it is hard to tell what is new etc!). Further, please do not
rebase your branch on main/force push/rewrite history, doing any of
these
   causes the context of any comments made by reviewers to be lost. If
   conflicts occur against main they should be resolved by merging main
   into the branch used for making the pull request.

Many thanks in advance for your cooperation!

-->
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[FEA] Add Numba's GDB pretty-printer to the Numba CUDA repository

2 participants