This is an archive of the discontinued LLVM Phabricator instance.

Differential D23460

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

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

Details

Reviewers
jholewinski
jlebar
Commits
rG2f0a3dfe64ab: [NVPTX] Use untyped (.b) integer registers in PTX.
rL278568: [NVPTX] Use untyped (.b) integer registers in PTX.
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.
Herald added a subscriber: jholewinski. · View Herald TranscriptAug 12 2016, 11:31 AM
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.

jlebar added a comment.Aug 12 2016, 12:14 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?

tra added a comment.Aug 12 2016, 12:32 PM
In D23460#514147, @jlebar wrote:

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.

jlebar added a comment.Aug 12 2016, 1:02 PM

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
Closed by commit rL278568: [NVPTX] Use untyped (.b) integer registers in PTX. (authored by tra). · Explain WhyAug 12 2016, 3:10 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
lib/
Target/
NVPTX/
6 lines
test/
CodeGen/
NVPTX/
69 lines

Diff 67883

lib/Target/NVPTX/NVPTXRegisterInfo.cpp

Show All 27 Lines
namespace llvm { namespace llvm {
std::string getNVPTXRegClassName(TargetRegisterClass const *RC) { std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
if (RC == &NVPTX::Float32RegsRegClass) { if (RC == &NVPTX::Float32RegsRegClass) {
return ".f32"; return ".f32";
} }
if (RC == &NVPTX::Float64RegsRegClass) { if (RC == &NVPTX::Float64RegsRegClass) {
return ".f64"; return ".f64";
} else if (RC == &NVPTX::Int64RegsRegClass) { } else if (RC == &NVPTX::Int64RegsRegClass) {
return ".s64"; return ".b64";
} else if (RC == &NVPTX::Int32RegsRegClass) { } else if (RC == &NVPTX::Int32RegsRegClass) {
return ".s32"; return ".b32";
jlebarUnsubmitted
Not Done
ReplyInline Actions

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.

jlebar: As written, this comment says we're making this change to avoid the hypothetical possibility of…
} else if (RC == &NVPTX::Int16RegsRegClass) { } else if (RC == &NVPTX::Int16RegsRegClass) {
return ".s16"; return ".b16";
} else if (RC == &NVPTX::Int1RegsRegClass) { } else if (RC == &NVPTX::Int1RegsRegClass) {
return ".pred"; return ".pred";
} else if (RC == &NVPTX::SpecialRegsRegClass) { } else if (RC == &NVPTX::SpecialRegsRegClass) {
return "!Special!"; return "!Special!";
} else { } else {
return "INTERNAL"; return "INTERNAL";
} }
return ""; return "";
▲ Show 20 Lines • Show All 62 Lines • Show Last 20 Lines

test/CodeGen/NVPTX/reg-types.ll

  • This file was added.
; Verify register types we generate in PTX.
; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT
; CHECK-LABEL: .visible .func func()
; NO8BIT-LABEL: .visible .func func()
define void @func() {
entry:
%s8 = alloca i8, align 1
%u8 = alloca i8, align 1
%s16 = alloca i16, align 2
%u16 = alloca i16, align 2
; Both 8- and 16-bit integers are packed into 16-bit registers.
; CHECK-DAG: .reg .b16 %rs<
; We should not generate 8-bit registers.
; NO8BIT-NOT: .reg .{{[bsu]}}8
%s32 = alloca i32, align 4
%u32 = alloca i32, align 4
; CHECK-DAG: .reg .b32 %r<
%s64 = alloca i64, align 8
%u64 = alloca i64, align 8
; CHECK-DAG: .reg .b64 %rd<
%f32 = alloca float, align 4
; CHECK-DAG: .reg .f32 %f<
jlebarUnsubmitted
Not Done
ReplyInline Actions

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

jlebar: I don't know what you're trying to check, but you can't do CHECK-NOT?
traAuthorUnsubmitted
Not Done
ReplyInline Actions

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.

tra: I need to make sure we did not generate `.reg .b8` for a register we use to store something in…
%f64 = alloca double, align 8
; CHECK-DAG: .reg .f64 %fd<
; Verify that we use correct register types.
store i8 1, i8* %s8, align 1
; CHECK: mov.u16 [[R1:%rs[0-9]]], 1;
; CHECK-NEXT: st.u8 {{.*}}, [[R1]]
store i8 2, i8* %u8, align 1
; CHECK: mov.u16 [[R2:%rs[0-9]]], 2;
; CHECK-NEXT: st.u8 {{.*}}, [[R2]]
store i16 3, i16* %s16, align 2
; CHECK: mov.u16 [[R3:%rs[0-9]]], 3;
; CHECK-NEXT: st.u16 {{.*}}, [[R3]]
store i16 4, i16* %u16, align 2
; CHECK: mov.u16 [[R4:%rs[0-9]]], 4;
; CHECK-NEXT: st.u16 {{.*}}, [[R4]]
store i32 5, i32* %s32, align 4
; CHECK: mov.u32 [[R5:%r[0-9]]], 5;
; CHECK-NEXT: st.u32 {{.*}}, [[R5]]
store i32 6, i32* %u32, align 4
; CHECK: mov.u32 [[R6:%r[0-9]]], 6;
; CHECK-NEXT: st.u32 {{.*}}, [[R6]]
store i64 7, i64* %s64, align 8
; CHECK: mov.u64 [[R7:%rd[0-9]]], 7;
; CHECK-NEXT: st.u64 {{.*}}, [[R7]]
store i64 8, i64* %u64, align 8
; CHECK: mov.u64 [[R8:%rd[0-9]]], 8;
; CHECK-NEXT: st.u64 {{.*}}, [[R8]]
; FP constants are stored via integer registers, but that's an
; implementation detail that's irrelevant here.
store float 9.000000e+00, float* %f32, align 4
store double 1.000000e+01, double* %f64, align 8
; Instead, we force a load into a register and then verify register type.
%f32v = load volatile float, float* %f32, align 4
; CHECK: ld.volatile.f32 %f{{[0-9]+}}
%f64v = load volatile double, double* %f64, align 8
; CHECK: ld.volatile.f64 %fd{{[0-9]+}}
ret void
; CHECK: ret;
; NO8BIT: ret;
}