From 46ef3dcd7e1598482e4d86688ef364fb4d5e1dca Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Tue, 16 Aug 2022 14:23:55 +0000 Subject: [PATCH 1/9] add pretty-printers --- scripts/gdb-pretty-printers.py | 229 +++++++++++++++++++++++++++++++++ 1 file changed, 229 insertions(+) create mode 100644 scripts/gdb-pretty-printers.py diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py new file mode 100644 index 000000000..972b4438a --- /dev/null +++ b/scripts/gdb-pretty-printers.py @@ -0,0 +1,229 @@ +# Copyright (c) 2022, NVIDIA CORPORATION. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# + +import gdb +import sys + +if sys.version_info[0] > 2: + Iterator = object +else: + # "Polyfill" for Python2 Iterator interface + class Iterator: + def next(self): + return self.__next__() + + +class HostIterator(Iterator): + """Iterates over arrays in host memory""" + + 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 DeviceIterator(Iterator): + """Iterates over arrays in device memory by copying chunks into host memory""" + + 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 + # At most 1 MB or size, at least 1 + 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) + + +class ThrustVectorPrinter(gdb.printing.PrettyPrinter): + """Print a thrust::*_vector""" + + 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 DeviceIterator(self.pointer, self.size) + else: + return HostIterator(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 + + +class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): + """Print a rmm::device_uvector""" + + def __init__(self, val): + self.val = val + el_type = val.type.template_argument(0) + self.pointer = val['_storage']['_data'].cast(el_type.pointer()) + self.size = int(val['_storage']['_size']) // el_type.sizeof + self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof + + def children(self): + return DeviceIterator(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' + + +# Workaround to avoid using the pretty printer on things like std::vector::iterator +def is_template_type_not_alias(typename): + loc = typename.find('<') + if loc is None: + return False + depth = 0 + for char in typename[typename.find('<'):-1]: + if char == '<': + depth += 1 + if char == '>': + depth -= 1 + if depth == 0: + return False + return True + + +def template_match(typename, template_name): + return typename.startswith(template_name + "<") and typename.endswith(">") + + +def lookup_rmm_type(val): + if not str(val.type.unqualified()).startswith('rmm::'): + return None + suffix = str(val.type.unqualified())[5:] + if not is_template_type_not_alias(suffix): + return None + if template_match(suffix, 'device_uvector'): + return RmmDeviceUVectorPrinter(val) + return None + + +def lookup_thrust_type(val): + if not str(val.type.unqualified()).startswith('thrust::'): + return None + suffix = str(val.type.unqualified())[8:] + if not is_template_type_not_alias(suffix): + return None + if template_match(suffix, 'host_vector') or template_match(suffix, 'device_vector'): + return ThrustVectorPrinter(val) + elif int(gdb.VERSION.split(".")[0]) >= 10 and template_match(suffix, 'device_reference'): + return ThrustReferencePrinter(val) + return None + + +gdb.pretty_printers.append(lookup_rmm_type) +gdb.pretty_printers.append(lookup_thrust_type) From 15bedd1d5e662603f08303eefe04c00a874b3402 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Tue, 16 Aug 2022 14:37:49 +0000 Subject: [PATCH 2/9] wip --- scripts/gdb-pretty-printers.py | 85 ++-------------------------------- 1 file changed, 4 insertions(+), 81 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index 972b4438a..eee1b3322 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -13,16 +13,12 @@ # limitations under the License. # -import gdb import sys +import gdb + -if sys.version_info[0] > 2: - Iterator = object -else: - # "Polyfill" for Python2 Iterator interface - class Iterator: - def next(self): - return self.__next__() +if "@RMM_PRETTY_PRINTER@" != "1": + sys.exit("This file is only a template, use gdb-pretty-printer.py in the CMake build directory instead.") class HostIterator(Iterator): @@ -101,65 +97,6 @@ def __next__(self): return ('[%d]' % count, elt) -class ThrustVectorPrinter(gdb.printing.PrettyPrinter): - """Print a thrust::*_vector""" - - 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 DeviceIterator(self.pointer, self.size) - else: - return HostIterator(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 - - class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): """Print a rmm::device_uvector""" @@ -212,18 +149,4 @@ def lookup_rmm_type(val): return None -def lookup_thrust_type(val): - if not str(val.type.unqualified()).startswith('thrust::'): - return None - suffix = str(val.type.unqualified())[8:] - if not is_template_type_not_alias(suffix): - return None - if template_match(suffix, 'host_vector') or template_match(suffix, 'device_vector'): - return ThrustVectorPrinter(val) - elif int(gdb.VERSION.split(".")[0]) >= 10 and template_match(suffix, 'device_reference'): - return ThrustReferencePrinter(val) - return None - - gdb.pretty_printers.append(lookup_rmm_type) -gdb.pretty_printers.append(lookup_thrust_type) From 533774e2132adfd296c37e0a401743eabc0a4bfe Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 17 Aug 2022 16:14:01 +0000 Subject: [PATCH 3/9] pull in thrust pretty-printer sources --- CMakeLists.txt | 7 + scripts/gdb-pretty-printers.py | 260 ++++++++++++++++----------------- 2 files changed, 137 insertions(+), 130 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 944fdbca1..63ee34fdd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -106,6 +106,13 @@ endif() include(CPack) +# optionally assemble Thrust pretty-printers +if(Thrust_SOURCE_DIR) + file(READ ${Thrust_SOURCE_DIR}/scripts/gdb-pretty-printers.py THRUST_PRETTY_PRINTER_CONTENT) + file(READ ${PROJECT_SOURCE_DIR}/scripts/gdb-pretty-printers.py RMM_PRETTY_PRINTER_CONTENT) + file(WRITE ${PROJECT_BINARY_DIR}/scripts/gdb-pretty-printers.py "${THRUST_PRETTY_PRINTER_CONTENT}${RMM_PRETTY_PRINTER_CONTENT}") +endif() + # install export targets install(TARGETS rmm EXPORT rmm-exports) install(DIRECTORY include/rmm/ DESTINATION include/rmm) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index eee1b3322..fbcd126be 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -13,140 +13,140 @@ # limitations under the License. # -import sys import gdb -if "@RMM_PRETTY_PRINTER@" != "1": - sys.exit("This file is only a template, use gdb-pretty-printer.py in the CMake build directory instead.") - - -class HostIterator(Iterator): - """Iterates over arrays in host memory""" - - 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 DeviceIterator(Iterator): - """Iterates over arrays in device memory by copying chunks into host memory""" - - 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 - # At most 1 MB or size, at least 1 - 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) - - -class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): - """Print a rmm::device_uvector""" - - def __init__(self, val): - self.val = val - el_type = val.type.template_argument(0) - self.pointer = val['_storage']['_data'].cast(el_type.pointer()) - self.size = int(val['_storage']['_size']) // el_type.sizeof - self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof - - def children(self): - return DeviceIterator(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' - - -# Workaround to avoid using the pretty printer on things like std::vector::iterator -def is_template_type_not_alias(typename): - loc = typename.find('<') - if loc is None: - return False - depth = 0 - for char in typename[typename.find('<'):-1]: - if char == '<': - depth += 1 - if char == '>': - depth -= 1 - if depth == 0: +if not 'ThrustVectorPrinter' in dir(): + print("This file expects the Thrust pretty-printers to be loaded already. " + "Either load them manually, or use the generated gdb-pretty-printers.py " + "in the build directory") +else: + class HostIterator(Iterator): + """Iterates over arrays in host memory""" + + 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 DeviceIterator(Iterator): + """Iterates over arrays in device memory by copying chunks into host memory""" + + 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 + # At most 1 MB or size, at least 1 + 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) + + + class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): + """Print a rmm::device_uvector""" + + def __init__(self, val): + self.val = val + el_type = val.type.template_argument(0) + self.pointer = val['_storage']['_data'].cast(el_type.pointer()) + self.size = int(val['_storage']['_size']) // el_type.sizeof + self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof + + def children(self): + return DeviceIterator(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' + + + # Workaround to avoid using the pretty printer on things like std::vector::iterator + def is_template_type_not_alias(typename): + loc = typename.find('<') + if loc is None: return False - return True - - -def template_match(typename, template_name): - return typename.startswith(template_name + "<") and typename.endswith(">") - - -def lookup_rmm_type(val): - if not str(val.type.unqualified()).startswith('rmm::'): - return None - suffix = str(val.type.unqualified())[5:] - if not is_template_type_not_alias(suffix): + depth = 0 + for char in typename[typename.find('<'):-1]: + if char == '<': + depth += 1 + if char == '>': + depth -= 1 + if depth == 0: + return False + return True + + + def template_match(typename, template_name): + return typename.startswith(template_name + "<") and typename.endswith(">") + + + def lookup_rmm_type(val): + if not str(val.type.unqualified()).startswith('rmm::'): + return None + suffix = str(val.type.unqualified())[5:] + if not is_template_type_not_alias(suffix): + return None + if template_match(suffix, 'device_uvector'): + return RmmDeviceUVectorPrinter(val) return None - if template_match(suffix, 'device_uvector'): - return RmmDeviceUVectorPrinter(val) - return None -gdb.pretty_printers.append(lookup_rmm_type) + gdb.pretty_printers.append(lookup_rmm_type) From 69eb5d27620c82c2a91a05efce8cd9b50cbfe2af Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 17 Aug 2022 18:20:30 +0000 Subject: [PATCH 4/9] provide script to load pretty-printers --- CMakeLists.txt | 4 +- scripts/gdb-pretty-printers.py | 261 ++++++++++++++++---------------- scripts/load-pretty-printers.in | 2 + 3 files changed, 134 insertions(+), 133 deletions(-) create mode 100644 scripts/load-pretty-printers.in diff --git a/CMakeLists.txt b/CMakeLists.txt index 63ee34fdd..9d81272ce 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -108,9 +108,7 @@ include(CPack) # optionally assemble Thrust pretty-printers if(Thrust_SOURCE_DIR) - file(READ ${Thrust_SOURCE_DIR}/scripts/gdb-pretty-printers.py THRUST_PRETTY_PRINTER_CONTENT) - file(READ ${PROJECT_SOURCE_DIR}/scripts/gdb-pretty-printers.py RMM_PRETTY_PRINTER_CONTENT) - file(WRITE ${PROJECT_BINARY_DIR}/scripts/gdb-pretty-printers.py "${THRUST_PRETTY_PRINTER_CONTENT}${RMM_PRETTY_PRINTER_CONTENT}") + configure_file(scripts/load-pretty-printers.in load-pretty-printers @ONLY) endif() # install export targets diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index fbcd126be..a34be3f2c 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -16,137 +16,138 @@ import gdb -if not 'ThrustVectorPrinter' in dir(): - print("This file expects the Thrust pretty-printers to be loaded already. " - "Either load them manually, or use the generated gdb-pretty-printers.py " - "in the build directory") -else: - class HostIterator(Iterator): - """Iterates over arrays in host memory""" - - 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 DeviceIterator(Iterator): - """Iterates over arrays in device memory by copying chunks into host memory""" - - 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 - # At most 1 MB or size, at least 1 - 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) - - - class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): - """Print a rmm::device_uvector""" - - def __init__(self, val): - self.val = val - el_type = val.type.template_argument(0) - self.pointer = val['_storage']['_data'].cast(el_type.pointer()) - self.size = int(val['_storage']['_size']) // el_type.sizeof - self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof - - def children(self): - return DeviceIterator(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' - - - # Workaround to avoid using the pretty printer on things like std::vector::iterator - def is_template_type_not_alias(typename): - loc = typename.find('<') - if loc is None: +if not 'ThrustVectorPrinter' in locals(): + raise Exception("This file expects the Thrust pretty-printers to be loaded already. " + "Either load them manually, or use the generated gdb-pretty-printers.py " + "in the build directory") + + +class HostIterator(Iterator): + """Iterates over arrays in host memory""" + + 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 DeviceIterator(Iterator): + """Iterates over arrays in device memory by copying chunks into host memory""" + + 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 + # At most 1 MB or size, at least 1 + 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) + + +class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): + """Print a rmm::device_uvector""" + + def __init__(self, val): + self.val = val + el_type = val.type.template_argument(0) + self.pointer = val['_storage']['_data'].cast(el_type.pointer()) + self.size = int(val['_storage']['_size']) // el_type.sizeof + self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof + + def children(self): + return DeviceIterator(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' + + +# Workaround to avoid using the pretty printer on things like std::vector::iterator +def is_template_type_not_alias(typename): + loc = typename.find('<') + if loc is None: + return False + depth = 0 + for char in typename[typename.find('<'):-1]: + if char == '<': + depth += 1 + if char == '>': + depth -= 1 + if depth == 0: return False - depth = 0 - for char in typename[typename.find('<'):-1]: - if char == '<': - depth += 1 - if char == '>': - depth -= 1 - if depth == 0: - return False - return True - - - def template_match(typename, template_name): - return typename.startswith(template_name + "<") and typename.endswith(">") - - - def lookup_rmm_type(val): - if not str(val.type.unqualified()).startswith('rmm::'): - return None - suffix = str(val.type.unqualified())[5:] - if not is_template_type_not_alias(suffix): - return None - if template_match(suffix, 'device_uvector'): - return RmmDeviceUVectorPrinter(val) + return True + + +def template_match(typename, template_name): + return typename.startswith(template_name + "<") and typename.endswith(">") + + +def lookup_rmm_type(val): + if not str(val.type.unqualified()).startswith('rmm::'): + return None + suffix = str(val.type.unqualified())[5:] + if not is_template_type_not_alias(suffix): return None + if template_match(suffix, 'device_uvector'): + return RmmDeviceUVectorPrinter(val) + return None - gdb.pretty_printers.append(lookup_rmm_type) +gdb.pretty_printers.append(lookup_rmm_type) diff --git a/scripts/load-pretty-printers.in b/scripts/load-pretty-printers.in new file mode 100644 index 000000000..bd59968cc --- /dev/null +++ b/scripts/load-pretty-printers.in @@ -0,0 +1,2 @@ +source @Thrust_SOURCE_DIR@/scripts/gdb-pretty-printers.py +source @PROJECT_SOURCE_DIR@/scripts/gdb-pretty-printers.py From 7c839a3fd05a35c84ddd43ffbd6ee302e8f2268a Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 17 Aug 2022 18:24:08 +0000 Subject: [PATCH 5/9] review updates --- scripts/gdb-pretty-printers.py | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index a34be3f2c..ac6c35bc5 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -18,8 +18,8 @@ if not 'ThrustVectorPrinter' in locals(): raise Exception("This file expects the Thrust pretty-printers to be loaded already. " - "Either load them manually, or use the generated gdb-pretty-printers.py " - "in the build directory") + "Either load them manually, or use the generated load-pretty-printers " + "script in the build directory") class HostIterator(Iterator): @@ -38,8 +38,8 @@ def __next__(self): raise StopIteration elt = self.item.dereference() count = self.count - self.item = self.item + 1 - self.count = self.count + 1 + self.item += 1 + self.count += 1 return ('[%d]' % count, elt) @@ -90,11 +90,11 @@ def __next__(self): raise StopIteration self.update_buffer() elt = self.buffer_item.dereference() - self.buffer_item = self.buffer_item + 1 - self.buffer_count = self.buffer_count + 1 + self.buffer_item += 1 + self.buffer_count += 1 count = self.count - self.item = self.item + 1 - self.count = self.count + 1 + self.item += 1 + self.count += 1 return ('[%d]' % count, elt) From 4171b2f9772f49c42d6954569c5ea09a5238df32 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 17 Aug 2022 18:40:12 +0000 Subject: [PATCH 6/9] modernize python code --- scripts/gdb-pretty-printers.py | 22 ++++++++++------------ 1 file changed, 10 insertions(+), 12 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index ac6c35bc5..f71945c36 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -22,7 +22,7 @@ "script in the build directory") -class HostIterator(Iterator): +class HostIterator(): """Iterates over arrays in host memory""" def __init__(self, start, size): @@ -40,10 +40,10 @@ def __next__(self): count = self.count self.item += 1 self.count += 1 - return ('[%d]' % count, elt) + return (f"[{count}]", elt) -class DeviceIterator(Iterator): +class DeviceIterator(): """Iterates over arrays in device memory by copying chunks into host memory""" def __init__(self, start, size): @@ -56,8 +56,7 @@ def __init__(self, start, size): self.buffer_start = 0 # At most 1 MB or size, at least 1 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 = gdb.parse_and_eval(f"(void*)malloc({self.buffer_size * self.sizeof})") self.buffer.fetch_lazy() self.buffer_count = self.buffer_size self.update_buffer() @@ -73,14 +72,13 @@ def update_buffer(self): 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)) + f"(cudaError)cudaMemcpy({buffer_addr}, {device_addr}, {size}, cudaMemcpyDeviceToHost)") if status != 0: raise gdb.MemoryError( - 'memcpy from device failed: %s' % status) + f"memcpy from device failed: {status}") def __del__(self): - gdb.parse_and_eval('(void)free(%s)' % - hex(self.buffer)).fetch_lazy() + gdb.parse_and_eval(f"(void)free({hex(self.buffer)})").fetch_lazy() def __iter__(self): return self @@ -95,7 +93,7 @@ def __next__(self): count = self.count self.item += 1 self.count += 1 - return ('[%d]' % count, elt) + return (f"[{count}]", elt) class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): @@ -113,7 +111,7 @@ def children(self): def to_string(self): typename = str(self.val.type) - return ('%s of length %d, capacity %d' % (typename, self.size, self.capacity)) + return (f"{typename} of length {self.size}, capacity {self.capacity}") def display_hint(self): return 'array' @@ -125,7 +123,7 @@ def is_template_type_not_alias(typename): if loc is None: return False depth = 0 - for char in typename[typename.find('<'):-1]: + for char in typename[loc:-1]: if char == '<': depth += 1 if char == '>': From 4e73504be8c817d722c5e67345e8c8a69b3550e1 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 17 Aug 2022 18:53:32 +0000 Subject: [PATCH 7/9] consistent quotes --- scripts/gdb-pretty-printers.py | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index f71945c36..e2a5cce1b 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -16,7 +16,7 @@ import gdb -if not 'ThrustVectorPrinter' in locals(): +if not "ThrustVectorPrinter" in locals(): raise Exception("This file expects the Thrust pretty-printers to be loaded already. " "Either load them manually, or use the generated load-pretty-printers " "script in the build directory") @@ -102,9 +102,9 @@ class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): def __init__(self, val): self.val = val el_type = val.type.template_argument(0) - self.pointer = val['_storage']['_data'].cast(el_type.pointer()) - self.size = int(val['_storage']['_size']) // el_type.sizeof - self.capacity = int(val['_storage']['_capacity']) // el_type.sizeof + self.pointer = val["_storage"]["_data"].cast(el_type.pointer()) + self.size = int(val["_storage"]["_size"]) // el_type.sizeof + self.capacity = int(val["_storage"]["_capacity"]) // el_type.sizeof def children(self): return DeviceIterator(self.pointer, self.size) @@ -114,19 +114,19 @@ def to_string(self): return (f"{typename} of length {self.size}, capacity {self.capacity}") def display_hint(self): - return 'array' + return "array" # Workaround to avoid using the pretty printer on things like std::vector::iterator def is_template_type_not_alias(typename): - loc = typename.find('<') + loc = typename.find("<") if loc is None: return False depth = 0 for char in typename[loc:-1]: - if char == '<': + if char == "<": depth += 1 - if char == '>': + if char == ">": depth -= 1 if depth == 0: return False @@ -138,12 +138,12 @@ def template_match(typename, template_name): def lookup_rmm_type(val): - if not str(val.type.unqualified()).startswith('rmm::'): + if not str(val.type.unqualified()).startswith("rmm::"): return None suffix = str(val.type.unqualified())[5:] if not is_template_type_not_alias(suffix): return None - if template_match(suffix, 'device_uvector'): + if template_match(suffix, "device_uvector"): return RmmDeviceUVectorPrinter(val) return None From ccfb1127ce328d6504cced09042807fcc3efdaac Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Tue, 23 Aug 2022 10:46:06 +0000 Subject: [PATCH 8/9] review updates --- scripts/gdb-pretty-printers.py | 17 ++++++++--------- 1 file changed, 8 insertions(+), 9 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index e2a5cce1b..1b48eb00c 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -16,14 +16,14 @@ import gdb -if not "ThrustVectorPrinter" in locals(): - raise Exception("This file expects the Thrust pretty-printers to be loaded already. " +if "ThrustVectorPrinter" not in locals(): + raise NameError("This file expects the Thrust pretty-printers to be loaded already. " "Either load them manually, or use the generated load-pretty-printers " "script in the build directory") -class HostIterator(): - """Iterates over arrays in host memory""" +class HostIterator: + """Iterates over arrays in host memory.""" def __init__(self, start, size): self.item = start @@ -43,8 +43,8 @@ def __next__(self): return (f"[{count}]", elt) -class DeviceIterator(): - """Iterates over arrays in device memory by copying chunks into host memory""" +class DeviceIterator: + """Iterates over arrays in device memory by copying chunks into host memory.""" def __init__(self, start, size): self.exec = exec @@ -97,7 +97,7 @@ def __next__(self): class RmmDeviceUVectorPrinter(gdb.printing.PrettyPrinter): - """Print a rmm::device_uvector""" + """Print a rmm::device_uvector.""" def __init__(self, val): self.val = val @@ -110,8 +110,7 @@ def children(self): return DeviceIterator(self.pointer, self.size) def to_string(self): - typename = str(self.val.type) - return (f"{typename} of length {self.size}, capacity {self.capacity}") + return (f"{self.val.type} of length {self.size}, capacity {self.capacity}") def display_hint(self): return "array" From 433fda05dda083ce5133c7615f6d6ec823b74a0c Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Tue, 23 Aug 2022 15:16:45 +0000 Subject: [PATCH 9/9] remove unnecessary error message --- scripts/gdb-pretty-printers.py | 6 ------ 1 file changed, 6 deletions(-) diff --git a/scripts/gdb-pretty-printers.py b/scripts/gdb-pretty-printers.py index 1b48eb00c..df6a662ee 100644 --- a/scripts/gdb-pretty-printers.py +++ b/scripts/gdb-pretty-printers.py @@ -16,12 +16,6 @@ import gdb -if "ThrustVectorPrinter" not in locals(): - raise NameError("This file expects the Thrust pretty-printers to be loaded already. " - "Either load them manually, or use the generated load-pretty-printers " - "script in the build directory") - - class HostIterator: """Iterates over arrays in host memory."""