Provide GDB pretty-printers for Thrust containers and iterators

[harrism]

It would be wonderful if I could inspect the contents of thrust containers: host_vector and device_vector in GDB (and more importantly, in VSCode). GDB allows customizing this.

It would save a lot of time if I could inspect device vectors without having to bring them to the host (e.g. the pretty printer script would do that behind the scenes for me).

Bonus: we also have custom containers with interfaces that look a lot like device_vector -- it would be great if Thrust's pretty-printers could be used on those as well.

Bonus 2: Moreover, since fancy iterators are so common in our Thrust code (notably "counting transform" iterators), it could be extremely valuable to be able to introspect the values of iterators. This might require running device code in the general case, but for example, if I want to look at a range of iterators, the pretty-printer might be able to run a thrust::sequence on a small range and then copy the results to the host and print them. Not sure if this is fully supportable, just thought it is a potentially useful idea.

[alliepiper]

This isn't something we can prioritize right now, but I agree it would be a very nice contribution.

[harrism]

Thanks @allisonvacanti. BTW @brycelelbach asked me to file this. Would be a nice enhancement for the CUDA developer experience in general.

[alliepiper]

He mentioned last week he might work on this at some point. I agree, it'd be good to add eventually.

[h20190907]

@allisonvacanti I would like to work on this issue. Could you please assign this to me?

[alliepiper]

Absolutely 😃

[h20190907]

@allisonvacanti @harrism GDB is used to debug code and visualize the working of the program. Could you please help me with the exact changes I need to make for this enhancement? The requirement is not clear to me

[alliepiper]

I'm afraid we're fairly shorthanded right now. This is a good first issue for someone already familiar with GDB and pretty printers, but we don't have spare cycles right now to provide a ton of help here.

If you need this functionality for your projects, I can point you at some resources, but if you're just looking to get involved with an open-source project we can probably find a better issue to start out.

In general, GDB pretty-printers are used to format complex data structures as strings, e.g. https://undo.io/resources/gdb-watchpoint/here-quick-way-pretty-print-structures-gdb/ and for thrust, there are some classes that would benefit from having them, which are listed in the original issue post.

[tomilov]

It seems that for thrust::device_vector GDB pretty printers can make sense only for CPP/TBB/OMP backends. All except CUDA.

[harrism]

It seems that for thrust::device_vector GDB pretty printers can make sense only for CPP/TBB/OMP backends. All except CUDA.

Why? Can GDB pretty printers not copy data? What I envision is that the pretty printer would do a cudaMemcpy to a host array in order to pretty print it.

[tomilov]

I don't think pretty printers can import and call some non-trivial functions from image of executable. Maybe I am wrong.

[harrism]

Pretty printers can be written in Python. I think one would just need python bindings to cudaMemcpy that could be imported into the Python pretty printer code. That could be used to get the data from a device_vector into a host array, which could then be accessed from the pretty printer.

https://www.kurokatta.org/grumble/2018/05/gdb-pretty

[alliepiper]

Sounds like I misunderstood the complexity involved in getting this to work. Removing the good first issue tag.

[alliepiper]

@h20190907 are you still interested in working on this, or should I update the assignment?

[tomilov]

There is gdb.parse_and_eval. It have to be called when process is actually running and in case of cudaMemcpy it can cause unrecoverable deadlock, if there is some synchronization inside it.