This is an archive of the discontinued LLVM Phabricator instance.

[NVPTX] Use untyped (.b) integer registers in PTX.
ClosedPublic

Authored by tra on Aug 12 2016, 11:31 AM.

Details

Summary

This brings LLVM-generated PTX closer to what nvcc generates and avoids
triggering issues in ptxas.

For instance, ptxas does not accept .s16 (or .u16) registers as operands
for .fp16 instructions.

Diff Detail

Event Timeline

tra updated this revision to Diff 67869.Aug 12 2016, 11:31 AM
tra retitled this revision from to [NVPTX] Use untyped (.b) integer registers in PTX..
tra updated this object.
tra added reviewers: jholewinski, jlebar.
tra added a subscriber: llvm-commits.
jlebar edited edge metadata.Aug 12 2016, 11:55 AM

Since we're working around a ptxas bug, I feel like we should have a comment somewhere explaining ourselves.

test/CodeGen/NVPTX/reg-types.ll
26

I don't know what you're trying to check, but you can't do CHECK-NOT?

tra added inline comments.Aug 12 2016, 12:10 PM
test/CodeGen/NVPTX/reg-types.ll
26

I need to make sure we did not generate .reg .b8 for a register we use to store something in %u8/%s8.
Instead, we expect generate .reg .b16. However, .reg .b16 is also used for operations on 16-bit %s16/%u16, so checking for its existence does not rule out .reg .b8. CHECK-NOT would work, if order of .reg would be guaranteed, but it's not.

According to FileCheck manual, CHECK-DAG entries around CHECK-NOT cannot be reordered. Considering that .reg lines are not ordered, placing CHECK-NOT anywhere in-between CHECK-DAG lines would break the tests.

CHECK-NOT would work, if order of .reg would be guaranteed, but it's not.

I believe that different --check-prefixes are evaluated independently, so if you did FOO-NOT, that would effectively check that something never appears in the file. Would that work?

tra added a comment.Aug 12 2016, 12:32 PM

CHECK-NOT would work, if order of .reg would be guaranteed, but it's not.

I believe that different --check-prefixes are evaluated independently, so if you did FOO-NOT, that would effectively check that something never appears in the file. Would that work?

Empirical evidence suggests that -check-prefixes are *not* evaluated independently. As far as I can tell, it just creates set of labels to pay attention to and then all labels within the set are treated the same.

That said, I can always do negative check in a separate RUN line, so the idea is helpful.

tra updated this revision to Diff 67883.Aug 12 2016, 12:38 PM
tra edited edge metadata.

Added negative test for not generating .reg .b8.

Test looks good to me, but I do think we want a comment in the source explaining why we use b.

tra updated this revision to Diff 67910.Aug 12 2016, 2:15 PM

Added a comment.

jlebar added inline comments.Aug 12 2016, 2:20 PM
lib/Target/NVPTX/NVPTXRegisterInfo.cpp
38

As written, this comment says we're making this change to avoid the hypothetical possibility of hitting bugs in ptxas in the future. But that's not really true, right? My understanding is there's an explicit bug in ptxas with fp16 that we are working around.

We chose to use this big hammer to work around the ptxas bug for the reasons you explain here, and that's also an important part, but it's not the whole story.

Sorry to be nitpicky about this, but spelling it out (ideally with examples of good/bad PTX) is really important so that if someone in the future comes along and wants to change this, they have some chance of being able to evaluate whether they will regress something.

tra updated this revision to Diff 67919.Aug 12 2016, 2:53 PM

Elaborated on reasons for using .b w/ example.

jlebar accepted this revision.Aug 12 2016, 2:54 PM
jlebar edited edge metadata.

Perfect. Thanks. :)

This revision is now accepted and ready to land.Aug 12 2016, 2:54 PM
This revision was automatically updated to reflect the committed changes.