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

[tra]

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.

[jlebar]

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

25 ↗

(On Diff #67869)

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

[tra]

test/CodeGen/NVPTX/reg-types.ll

25 ↗

(On Diff #67869)

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]

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]

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]

Added negative test for not generating .reg .b8.

[jlebar]

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

[tra]

Added a comment.

[jlebar]

lib/Target/NVPTX/NVPTXRegisterInfo.cpp

38 ↗

(On Diff #67910)

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]

Elaborated on reasons for using .b w/ example.

[jlebar]

Perfect. Thanks. :)

llvm/trunk/lib/Target/NVPTX/NVPTXRegisterInfo.cpp

 namespace llvm {
 std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
   if (RC == &NVPTX::Float32RegsRegClass) {
     return ".f32";
   }
   if (RC == &NVPTX::Float64RegsRegClass) {
     return ".f64";
   } else if (RC == &NVPTX::Int64RegsRegClass) {
-    return ".s64";
+    // We use untyped (.b) integer registers here as NVCC does.
+    // Correctness of generated code does not depend on register type,
+    // but using .s/.u registers runs into ptxas bug that prevents
+    // assembly of otherwise valid PTX into SASS. Despite PTX ISA
+    // specifying only argument size for fp16 instructions, ptxas does
+    // not allow using .s16 or .u16 arguments for .fp16
+    // instructions. At the same time it allows using .s32/.u32
+    // arguments for .fp16v2 instructions:
+    //
+    //   .reg .b16 rb16
+    //   .reg .s16 rs16
+    //   add.f16 rb16,rb16,rb16; // OK
+    //   add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add'
+    // but:
+    //   .reg .b32 rb32
+    //   .reg .s32 rs32
+    //   add.f16v2 rb32,rb32,rb32; // OK
+    //   add.f16v2 rs32,rs32,rs32; // OK
+    return ".b64";
   } else if (RC == &NVPTX::Int32RegsRegClass) {
-    return ".s32";
+    return ".b32";
   } else if (RC == &NVPTX::Int16RegsRegClass) {
-    return ".s16";
+    return ".b16";
   } else if (RC == &NVPTX::Int1RegsRegClass) {
     return ".pred";
   } else if (RC == &NVPTX::SpecialRegsRegClass) {
     return "!Special!";
   } else {
     return "INTERNAL";
   }
   return "";

llvm/trunk/test/CodeGen/NVPTX/reg-types.ll

+; 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<
+  %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;
+}
+