Page MenuHomePhabricator

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

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



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
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

rC Clang

Event Timeline

ABataev created this revision.Jan 24 2019, 8:27 AM
probinson added inline comments.Jan 30 2019, 9:13 AM

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.

__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

This revision was automatically updated to reflect the committed changes.