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

This bring 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.

Differential Revision: https://reviews.llvm.org/D23460

llvm-svn: 278568

diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 6f8672f..6cbf060 100644 (file)
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -33,11 +33,29 @@ std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
   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) {

diff --git a/llvm/test/CodeGen/NVPTX/reg-types.ll b/llvm/test/CodeGen/NVPTX/reg-types.ll
new file mode 100644 (file)
index 0000000..25494fa
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/reg-types.ll
@@ -0,0 +1,69 @@
+; 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;
+}
+