| import gdb |
| import sys |
|
|
| if sys.version_info[0] > 2: |
| Iterator = object |
| else: |
| |
| class Iterator: |
| def next(self): |
| return self.__next__() |
|
|
|
|
| class ThrustVectorPrinter(gdb.printing.PrettyPrinter): |
| "Print a thrust::*_vector" |
|
|
| class _host_accessible_iterator(Iterator): |
| def __init__(self, start, size): |
| self.item = start |
| self.size = size |
| self.count = 0 |
|
|
| def __iter__(self): |
| return self |
|
|
| def __next__(self): |
| if self.count >= self.size: |
| raise StopIteration |
| elt = self.item.dereference() |
| count = self.count |
| self.item = self.item + 1 |
| self.count = self.count + 1 |
| return ('[%d]' % count, elt) |
|
|
| class _device_iterator(Iterator): |
| def __init__(self, start, size): |
| self.exec = exec |
| self.item = start |
| self.size = size |
| self.count = 0 |
| self.buffer = None |
| self.sizeof = self.item.dereference().type.sizeof |
| self.buffer_start = 0 |
| |
| self.buffer_size = min(size, max(1, 2 ** 20 // self.sizeof)) |
| self.buffer = gdb.parse_and_eval( |
| '(void*)malloc(%s)' % (self.buffer_size * self.sizeof)) |
| self.buffer.fetch_lazy() |
| self.buffer_count = self.buffer_size |
| self.update_buffer() |
|
|
| def update_buffer(self): |
| if self.buffer_count >= self.buffer_size: |
| self.buffer_item = gdb.parse_and_eval( |
| hex(self.buffer)).cast(self.item.type) |
| self.buffer_count = 0 |
| self.buffer_start = self.count |
| device_addr = hex(self.item.dereference().address) |
| buffer_addr = hex(self.buffer) |
| size = min(self.buffer_size, self.size - |
| self.buffer_start) * self.sizeof |
| status = gdb.parse_and_eval( |
| '(cudaError)cudaMemcpy(%s, %s, %d, cudaMemcpyDeviceToHost)' % (buffer_addr, device_addr, size)) |
| if status != 0: |
| raise gdb.MemoryError( |
| 'memcpy from device failed: %s' % status) |
|
|
| def __del__(self): |
| gdb.parse_and_eval('(void)free(%s)' % |
| hex(self.buffer)).fetch_lazy() |
|
|
| def __iter__(self): |
| return self |
|
|
| def __next__(self): |
| if self.count >= self.size: |
| raise StopIteration |
| self.update_buffer() |
| elt = self.buffer_item.dereference() |
| self.buffer_item = self.buffer_item + 1 |
| self.buffer_count = self.buffer_count + 1 |
| count = self.count |
| self.item = self.item + 1 |
| self.count = self.count + 1 |
| return ('[%d]' % count, elt) |
|
|
| def __init__(self, val): |
| self.val = val |
| self.pointer = val['m_storage']['m_begin']['m_iterator'] |
| self.size = int(val['m_size']) |
| self.capacity = int(val['m_storage']['m_size']) |
| self.is_device = False |
| if str(self.pointer.type).startswith("thrust::device_ptr"): |
| self.pointer = self.pointer['m_iterator'] |
| self.is_device = True |
|
|
| def children(self): |
| if self.is_device: |
| return self._device_iterator(self.pointer, self.size) |
| else: |
| return self._host_accessible_iterator(self.pointer, self.size) |
|
|
| def to_string(self): |
| typename = str(self.val.type) |
| return ('%s of length %d, capacity %d' % (typename, self.size, self.capacity)) |
|
|
| def display_hint(self): |
| return 'array' |
|
|
|
|
| class ThrustReferencePrinter(gdb.printing.PrettyPrinter): |
| "Print a thrust::device_reference" |
|
|
| def __init__(self, val): |
| self.val = val |
| self.pointer = val['ptr']['m_iterator'] |
| self.type = self.pointer.dereference().type |
| sizeof = self.type.sizeof |
| self.buffer = gdb.parse_and_eval('(void*)malloc(%s)' % sizeof) |
| device_addr = hex(self.pointer) |
| buffer_addr = hex(self.buffer) |
| status = gdb.parse_and_eval('(cudaError)cudaMemcpy(%s, %s, %d, cudaMemcpyDeviceToHost)' % ( |
| buffer_addr, device_addr, sizeof)) |
| if status != 0: |
| raise gdb.MemoryError('memcpy from device failed: %s' % status) |
| self.buffer_val = gdb.parse_and_eval( |
| hex(self.buffer)).cast(self.pointer.type).dereference() |
|
|
| def __del__(self): |
| gdb.parse_and_eval('(void)free(%s)' % hex(self.buffer)).fetch_lazy() |
|
|
| def children(self): |
| return [] |
|
|
| def to_string(self): |
| typename = str(self.val.type) |
| return ('(%s) @%s: %s' % (typename, self.pointer, self.buffer_val)) |
|
|
| def display_hint(self): |
| return None |
|
|
|
|
| def lookup_thrust_type(val): |
| if not str(val.type.unqualified()).startswith('thrust::'): |
| return None |
| suffix = str(val.type.unqualified())[8:] |
| if suffix.startswith('host_vector') or suffix.startswith('device_vector'): |
| return ThrustVectorPrinter(val) |
| elif int(gdb.VERSION.split(".")[0]) >= 10 and suffix.startswith('device_reference'): |
| return ThrustReferencePrinter(val) |
| return None |
|
|
|
|
| gdb.pretty_printers.append(lookup_thrust_type) |
|
|
