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

Repository
rC Clang

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.