diff --git a/sycl/gdb/libsycl.so-gdb.py b/sycl/gdb/libsycl.so-gdb.py index 91ef325387976..5346621d18c35 100644 --- a/sycl/gdb/libsycl.so-gdb.py +++ b/sycl/gdb/libsycl.so-gdb.py @@ -22,6 +22,20 @@ def strip_version(name): return name.replace("sycl::_V1::", "sycl::") +def linear_index(gdb_range, gdb_id): + """Calculate the linear index for a sycl id within the specified range.""" + sycl_range = SYCLRange(gdb_range) + sycl_id = SYCLId(gdb_id) + assert sycl_range.dimensions() == sycl_id.dimensions() + index_array = sycl_id.array() + range_array = sycl_range.array() + linear_index = 0 + for dim in range(sycl_range.dimensions()): + linear_index *= range_array[dim] + linear_index += index_array[dim] + return linear_index + + class SYCLType: """Wrapper around a gdb.Type.""" @@ -269,7 +283,11 @@ def data(self): class SYCLLocalAccessor(SYCLValue): - """Provides information about a sycl::local_accessor from a gdb.Value.""" + """Provides information about a sycl::local_accessor from a gdb.Value. + + This SYCLLocalAccessor cannot represent the elements as a vector because + the address space qualifier only applies to pointers. + """ def __init__(self, gdb_value): super().__init__(gdb_value) @@ -304,15 +322,14 @@ def data(self): return data def subscript_sizet(self, subscript): - return self.data()[subscript] + if self.dimensions() == 1: + return self.data_ptr()[subscript] + else: + return self.data()[subscript] def subscript_id(self, subscript): - id_array = SYCLId(subscript).array() - result = self.data() - for dimension in range(self.dimensions()): - index = id_array[dimension] - result = result[index] - return result + index = linear_index(self.memory_range(), subscript) + return self.data_ptr()[index] def subscript_item(self, subscript): return self.subscript_id(SYCLItem(subscript).index()) @@ -757,6 +774,10 @@ def children(self): class SYCLLocalAccessorPrinter(SYCLPrinter): """Pretty printer for a sycl::local_accessor. + Use a decorated pointer for printing data in local memory. This implements + a trivial array printing algorithm by calculating indexes from a decorated + pointer. + Examples: # In device code. (gdb) p local_accessor_1d @@ -765,13 +786,63 @@ class SYCLLocalAccessorPrinter(SYCLPrinter): # In host code. (gdb) p local_accessor_1d $2 = sycl::local_accessor undefined + + In order to avoid printing arbitrarily large arrays, this printer includes + support for specifying the maximum elements. For example: + + Example: + (gdb) show print elements + Limit on string chars or array elements to print is 200. + (gdb) p -elements 4 -- local_accessor_1d + $3 = sycl::local_accessor range 10 = {0, 1, 2, 3...} """ def __init__(self, gdb_value): super().__init__(gdb_value) - def display_hint(self): - return "array" + def elements_per_dimension(self, gdb_range): + sycl_range = SYCLRange(gdb_range) + dims = sycl_range.dimensions() + extents = [int(sycl_range.array()[i]) for i in range(dims)] + product = lambda x: x[0] if len(x) == 1 else x[0] * product(x[1:]) + elements = [product(extents[dim:]) for dim in range(dims)] + return elements + + def array_string(self, gdb_pointer, gdb_range): + # Maximum number of elements for a single array dimension. + # print_options() is only available in GDB versions after 13.1. + if hasattr(gdb, "print_options") and callable(gdb.print_options): + max_elements = int(gdb.print_options()["max_elements"]) + else: + max_elements = int(gdb.parameter("print elements")) + string = "" + EPD = self.elements_per_dimension(gdb_range) + num_elements = EPD[0] + # Iterate over every item in the mulit-dimensional array. + item = 0 + while item < num_elements: + # Seperate items with a comma. + if item != 0: + string += ", " + for elements in EPD: + # Begin each array with '{'. + if item % elements == 0: + string += "{" + # Add the current array element. + string += str(gdb_pointer[item]) + sub_elements = 1 + for elements in reversed(EPD): + # Upon reaching max_elements in a single array, display "..." + # and skip the remaining elements in this array. + if ((item + 1) % elements) // sub_elements == max_elements: + string += "..." + item += elements - ((item + 1) % elements) + # End each array with '}'. + if (item + 1) % elements == 0: + string += "}" + sub_elements = elements + item += 1 + return string def to_string(self): local_accessor = SYCLLocalAccessor(self.gdb_value()) @@ -780,21 +851,13 @@ def to_string(self): mem_range = local_accessor.memory_range() range_string = SYCLRangePrinter(mem_range).value_as_string() string += " range " + range_string - except: + pointer = local_accessor.data_ptr() + array_string = self.array_string(pointer, mem_range) + string += " = " + array_string + except Exception: string += " undefined" return string - def children(self): - try: - local_accessor = SYCLLocalAccessor(self.gdb_value()) - data = local_accessor.data() - (low, inclusive_high) = data.type.range() - high = inclusive_high + 1 - for index in range(low, high): - yield (f"[{index}]", data[index]) - except: - return None - class SYCLBufferPrinter(SYCLPrinter): """Pretty printer for a sycl::buffer.