Page MenuHomePhabricator

[clang] Use correct address space for global variable debug info

Authored by scott.linder on Oct 7 2020, 9:07 AM.



The target needs to be queried here, but previously we seemed to only
duplicate CUDA's (and so HIP's) behavior, and only partially. Use the
same function as codegen to find the correct address space.

Diff Detail

Event Timeline

scott.linder created this revision.Oct 7 2020, 9:07 AM
Herald added a project: Restricted Project. · View Herald TranscriptOct 7 2020, 9:07 AM
scott.linder requested review of this revision.Oct 7 2020, 9:07 AM
scott.linder added a comment.EditedOct 7 2020, 9:14 AM

I'm not certain I fully understand NVPTX's relationship with its debugger, but from I gather the "default" address space in the debugger is global, and so the frontend omits it rather than explicitly mentioning it. I think it would be simpler to carry this information throughout the compiler, and only strip it late in the backend as a quirk controllable via some "optimize for NVPTX debugger", but in the patch as it currently is I instead just update NVPTXDWARFAddrSpaceMap.

Edit: Concerning auto variables, when coming back to the patch to post it I had missed the next patch in the series which addresses it by directly referring to the corresponding alloca rather than the addrspacecast to the default address space. I'll post that patch shortly to address the "FIXME" in this one.

aprantl accepted this revision.Oct 9 2020, 11:54 AM

That looks much nicer.


They are more convenient, but having very many CHECK_DAGs is also really slow — would it be feasible to reorder them and use CHECKs? Perhaps by running clang/FileCheck twice with different sets of CHECK lines?

This revision is now accepted and ready to land.Oct 9 2020, 11:54 AM

Replace uses of CHECK-DAG, use more meaningful names in test

scott.linder added inline comments.Oct 9 2020, 12:35 PM

The test is pretty short, so I just re-ordered the checks to match how they appear in the output (and used more descriptive names to make it easier to follow).

This does mean the test relies on the order these things are traversed. Some bits are maybe a bit surprising, like how the metadata for the __shared__ auto variable comes before the argument, but I don't imagine it is liable to change often/accidentally.

scott.linder added inline comments.Oct 14 2020, 12:10 PM

Does anyone have any thoughts on this change specifically? Is someone more familiar with NVPTX willing to weigh in on whether it makes more sense to carry the address space throughout the compiler explicitly and "drop" it late in the DWARF emission, or to do what I did in the current patch (drop it early).

I would lean towards updating the patch to do the latter, but I wanted to get feedback before plunging off to do it.

Herald added a project: Restricted Project. · View Herald TranscriptSep 28 2022, 2:27 PM
yaxunl accepted this revision.Sep 29 2022, 8:18 AM

LGTM. Thanks.