Page MenuHomePhabricator

[DEBUG_INFO][NVPTX] Generate correct data about variable address class.
ClosedPublic

Authored by ABataev on Jan 24 2019, 8:27 AM.

Details

Summary

Added ability to generate correct debug info data about the variable
address class. Currently, for all the locals and globals the default
values are used, ADDR_local_space(6) for locals and ADDR_global_space(5)
for globals. The values are taken from the table in

https://docs.nvidia.com/cuda/archive/10.0/ptx-writers-guide-to-interoperability/index.html#cuda-specific-dwarf.
We need to emit correct data for address classes of, at least, shared
and constant globals. Currently, all these variables are treated by
the cuda-gdb debugger as the variables in the global address space
and, thus, it require manual data type casting.

Diff Detail

Event Timeline

ABataev created this revision.Jan 24 2019, 8:27 AM
probinson added inline comments.Jan 30 2019, 9:13 AM
lib/CodeGen/CGDebugInfo.cpp
4235

Can a variable have one of these CUDA attributes when CUDAIsDevice is false? I'm just wondering if the extra level of checking is really necessary or useful.

tra added a subscriber: tra.Jan 30 2019, 9:57 AM
tra added inline comments.
lib/CodeGen/CGDebugInfo.cpp
4235

__shared__ are probably not going to be encountered on the host side, but __constant__ ones may be. E.g. we may have a global __constant__ variable which will have a host-side address (it's actually the address of its host-side 'shadow') which may be used by various CUDA functions to refer to it's real device-side instance.

probinson accepted this revision.Feb 1 2019, 7:50 AM

LGTM. I'll trust you on the actual address-class values.

This revision is now accepted and ready to land.Feb 1 2019, 7:50 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 1 2019, 7:50 AM

LGTM. I'll trust you on the actual address-class values.

Thanks, Paul! If you want, you can find those values in the table here https://docs.nvidia.com/cuda/archive/10.0/ptx-writers-guide-to-interoperability/index.html#cuda-specific-dwarf.

This revision was automatically updated to reflect the committed changes.