From ed43650bc22bfbd0bd04d6293de3e17f75222a13 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Sat, 27 Aug 2022 00:49:15 +0200 Subject: [PATCH] Add `gdb` pretty-printers for rmm types (#1088) This PR adds a pretty-printer for `device_uvector` and pulls in Thrust pretty-printers that were added in https://github.com/NVIDIA/thrust/pull/1631. CMake provides a convenience script to load all of the pretty-printers, to resolve the duplication concerns raised in https://github.com/rapidsai/cudf/pull/11499. Example output:
``` $ gdb -q gtests/DEVICE_UVECTOR_TEST Reading symbols from gtests/DEVICE_UVECTOR_TEST... (gdb) b cudaMalloc Function "cudaMalloc" not defined. Make breakpoint pending on future shared library load? (y or [n]) y Breakpoint 1 (cudaMalloc) pending. (gdb) run Starting program: /home/nfs/tribizel/rapids/rmm/build/cuda-11.5.0/feature__pretty-printers/debug/gtests/DEVICE_UVECTOR_TEST warning: Error disabling address space randomization: Operation not permitted [Thread debugging using libthread_db enabled] Using host libthread_db library "/usr/lib/x86_64-linux-gnu/libthread_db.so.1". Running main() from gmock_main.cc [==========] Running 95 tests from 5 test suites. [----------] Global test environment set-up. [----------] 19 tests from TypedUVectorTest/0, where TypeParam = signed char [ RUN ] TypedUVectorTest/0.MemoryResource [New Thread 0x7f741efc1000 (LWP 86147)] Thread 1 "DEVICE_UVECTOR_" hit Breakpoint 1, 0x00007f7426dddc80 in cudaMalloc () from /home/nfs/tribizel/rapids/compose/etc/conda/cuda_11.5/envs/rapids/lib/libcudart.so.11.0 (gdb) c Continuing. [New Thread 0x7f741e7c0000 (LWP 86148)] [ OK ] TypedUVectorTest/0.MemoryResource (999 ms) [ RUN ] TypedUVectorTest/0.ZeroSizeConstructor [ OK ] TypedUVectorTest/0.ZeroSizeConstructor (0 ms) [ RUN ] TypedUVectorTest/0.NonZeroSizeConstructor Thread 1 "DEVICE_UVECTOR_" hit Breakpoint 1, 0x00007f7426dddc80 in cudaMalloc () from /home/nfs/tribizel/rapids/compose/etc/conda/cuda_11.5/envs/rapids/lib/libcudart.so.11.0 (gdb) finish Run till exit from #0 0x00007f7426dddc80 in cudaMalloc () from /home/nfs/tribizel/rapids/compose/etc/conda/cuda_11.5/envs/rapids/lib/libcudart.so.11.0 rmm::mr::cuda_memory_resource::do_allocate (bytes=, this=) at /home/nfs/tribizel/rapids/rmm/include/rmm/mr/device/cuda_memory_resource.hpp:70 70 RMM_CUDA_TRY_ALLOC(cudaMalloc(&ptr, bytes)); (gdb) finish Run till exit from #0 rmm::mr::cuda_memory_resource::do_allocate (bytes=, this=) at /home/nfs/tribizel/rapids/rmm/include/rmm/mr/device/cuda_memory_resource.hpp:70 0x00005587874feef0 in rmm::device_buffer::allocate_async (bytes=12345, this=0x7ffefebb46b0) at /home/nfs/tribizel/rapids/rmm/include/rmm/device_buffer.hpp:418 418 _data = (bytes > 0) ? memory_resource()->allocate(bytes, stream()) : nullptr; Value returned is $1 = (void *) 0x7f73ff000000 (gdb) finish Run till exit from #0 0x00005587874feef0 in rmm::device_buffer::allocate_async (bytes=12345, this=0x7ffefebb46b0) at /home/nfs/tribizel/rapids/rmm/include/rmm/device_buffer.hpp:418 TypedUVectorTest_NonZeroSizeConstructor_Test::TestBody (this=) at /home/nfs/tribizel/rapids/rmm/tests/device_uvector_tests.cpp:55 55 EXPECT_EQ(vec.size(), size); (gdb) print vec $2 = {_storage = {_data = 0x7f73ff000000, _size = 12345, _capacity = 12345, _stream = {stream_ = 0x0}, _mr = 0x558787561008 }} (gdb) source load-pretty-printers (gdb) print vec $3 = rmm::device_uvector of length 12345, capacity 12345 = {0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000', 0 '\000'...} (gdb) ```
Authors: - Tobias Ribizel (https://github.com/upsj) Approvers: - Mark Harris (https://github.com/harrism) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/rmm/pull/1088 --- CMakeLists.txt | 5 ++ scripts/gdb-pretty-printers.py | 144 ++++++++++++++++++++++++++++++++ scripts/load-pretty-printers.in | 2 + 3 files changed, 151 insertions(+) create mode 100644 scripts/gdb-pretty-printers.py create mode 100644 scripts/load-pretty-printers.in diff --git a/CMakeLists.txt b/CMakeLists.txt index ce246596e..d5de4f6c7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -108,6 +108,11 @@ endif() include(CPack) +# optionally assemble Thrust pretty-printers +if(Thrust_SOURCE_DIR) + configure_file(scripts/load-pretty-printers.in load-pretty-printers @ONLY) +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 new file mode 100644 index 000000000..df6a662ee --- /dev/null +++ b/scripts/gdb-pretty-printers.py @@ -0,0 +1,144 @@ +# 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 + + +class HostIterator: + """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 += 1 + self.count += 1 + return (f"[{count}]", elt) + + +class DeviceIterator: + """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(f"(void*)malloc({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( + f"(cudaError)cudaMemcpy({buffer_addr}, {device_addr}, {size}, cudaMemcpyDeviceToHost)") + if status != 0: + raise gdb.MemoryError( + f"memcpy from device failed: {status}") + + def __del__(self): + gdb.parse_and_eval(f"(void)free({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 += 1 + self.buffer_count += 1 + count = self.count + self.item += 1 + self.count += 1 + return (f"[{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): + return (f"{self.val.type} of length {self.size}, capacity {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[loc:-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 + + +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