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 NVIDIA/thrust#1631. CMake provides a convenience script to load all of the pretty-printers, to resolve the duplication concerns raised in rapidsai/cudf#11499.

Example output:
<details>

```
$ 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=<optimized out>, this=<optimized out>) 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=<optimized out>, this=<optimized out>) 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<signed char>::TestBody (this=<optimized out>) 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 <rmm::mr::detail::initial_resource()::mr>}}
(gdb) source load-pretty-printers 
(gdb) print vec
$3 = rmm::device_uvector<signed char> 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) 
```

</details>

Authors:
  - Tobias Ribizel (https://github.com/upsj)

Approvers:
  - Mark Harris (https://github.com/harrism)
  - Bradley Dice (https://github.com/bdice)

URL: #1088

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)

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<int>::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)

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