[NVPTX] Switch to untyped float registers (#137011)
Register types in PTX are simply syntactic sugar and emitting them has
added lots of unnecessary complexity to the NVPTX backend. This change
takes the first step to their removal by using ".b" registers instead of
".f" in all cases. This should shake out any potential issues or bugs in
ptxas preventing full removal and pre-fetches many of the required test
updates.
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 1423888..6b9797c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -26,9 +26,9 @@
namespace llvm {
StringRef getNVPTXRegClassName(TargetRegisterClass const *RC) {
if (RC == &NVPTX::Float32RegsRegClass)
- return ".f32";
+ return ".b32";
if (RC == &NVPTX::Float64RegsRegClass)
- return ".f64";
+ return ".b64";
if (RC == &NVPTX::Int128RegsRegClass)
return ".b128";
if (RC == &NVPTX::Int64RegsRegClass)
diff --git a/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll b/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
index e46657e..8f0964c 100644
--- a/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
+++ b/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
@@ -45,7 +45,7 @@
; ENABLED-LABEL: fh(
; ENABLED: {
; ENABLED-NEXT: .reg .b16 %rs<10>;
-; ENABLED-NEXT: .reg .f32 %f<13>;
+; ENABLED-NEXT: .reg .b32 %f<13>;
; ENABLED-NEXT: .reg .b64 %rd<2>;
; ENABLED-EMPTY:
; ENABLED-NEXT: // %bb.0:
@@ -74,7 +74,7 @@
; DISABLED-LABEL: fh(
; DISABLED: {
; DISABLED-NEXT: .reg .b16 %rs<10>;
-; DISABLED-NEXT: .reg .f32 %f<13>;
+; DISABLED-NEXT: .reg .b32 %f<13>;
; DISABLED-NEXT: .reg .b64 %rd<2>;
; DISABLED-EMPTY:
; DISABLED-NEXT: // %bb.0:
@@ -121,7 +121,7 @@
define float @ff(ptr %p) {
; ENABLED-LABEL: ff(
; ENABLED: {
-; ENABLED-NEXT: .reg .f32 %f<10>;
+; ENABLED-NEXT: .reg .b32 %f<10>;
; ENABLED-NEXT: .reg .b64 %rd<2>;
; ENABLED-EMPTY:
; ENABLED-NEXT: // %bb.0:
@@ -137,7 +137,7 @@
;
; DISABLED-LABEL: ff(
; DISABLED: {
-; DISABLED-NEXT: .reg .f32 %f<10>;
+; DISABLED-NEXT: .reg .b32 %f<10>;
; DISABLED-NEXT: .reg .b64 %rd<2>;
; DISABLED-EMPTY:
; DISABLED-NEXT: // %bb.0:
diff --git a/llvm/test/CodeGen/NVPTX/and-or-setcc.ll b/llvm/test/CodeGen/NVPTX/and-or-setcc.ll
index 6c3514c..5949de3 100644
--- a/llvm/test/CodeGen/NVPTX/and-or-setcc.ll
+++ b/llvm/test/CodeGen/NVPTX/and-or-setcc.ll
@@ -9,7 +9,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [and_ord_param_0];
@@ -29,7 +29,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [or_uno_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/atomics.ll b/llvm/test/CodeGen/NVPTX/atomics.ll
index bb04aa8..16de80d 100644
--- a/llvm/test/CodeGen/NVPTX/atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics.ll
@@ -351,7 +351,7 @@
define float @atomic_add_f32_generic(ptr %addr, float %val) {
; CHECK-LABEL: atomic_add_f32_generic(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -370,7 +370,7 @@
define float @atomic_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
; CHECK-LABEL: atomic_add_f32_addrspace1(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -389,7 +389,7 @@
define float @atomic_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {
; CHECK-LABEL: atomic_add_f32_addrspace3(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -406,7 +406,7 @@
define float @atomicrmw_add_f32_generic(ptr %addr, float %val) {
; CHECK-LABEL: atomicrmw_add_f32_generic(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -426,7 +426,7 @@
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<17>;
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -470,7 +470,7 @@
define float @atomicrmw_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
; CHECK-LABEL: atomicrmw_add_f32_addrspace1(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -487,7 +487,7 @@
define float @atomicrmw_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {
; CHECK-LABEL: atomicrmw_add_f32_addrspace3(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
diff --git a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
index b97cb6f..6be13c3 100644
--- a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
@@ -19,7 +19,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<11>;
-; SM70-NEXT: .reg .f32 %f<4>;
+; SM70-NEXT: .reg .b32 %f<4>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fadd_param_1];
@@ -55,7 +55,7 @@
; SM80-FTZ-LABEL: test_fadd(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<4>;
+; SM80-FTZ-NEXT: .reg .b32 %f<4>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];
@@ -87,7 +87,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<11>;
-; SM70-NEXT: .reg .f32 %f<4>;
+; SM70-NEXT: .reg .b32 %f<4>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fsub_param_1];
@@ -123,7 +123,7 @@
; SM80-FTZ-LABEL: test_fsub(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<4>;
+; SM80-FTZ-NEXT: .reg .b32 %f<4>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];
@@ -155,7 +155,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<5>;
; SM70-NEXT: .reg .b32 %r<24>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];
@@ -210,7 +210,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-NEXT: .reg .b32 %f<7>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];
@@ -247,7 +247,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<5>;
; SM70-NEXT: .reg .b32 %r<24>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];
@@ -302,7 +302,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-NEXT: .reg .b32 %f<7>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];
@@ -339,7 +339,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<5>;
; SM70-NEXT: .reg .b32 %r<24>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];
@@ -394,7 +394,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-NEXT: .reg .b32 %f<7>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];
@@ -431,7 +431,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<5>;
; SM70-NEXT: .reg .b32 %r<24>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
@@ -474,7 +474,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<5>;
; SM80-NEXT: .reg .b32 %r<4>;
-; SM80-NEXT: .reg .f32 %f<7>;
+; SM80-NEXT: .reg .b32 %f<7>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
@@ -495,7 +495,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
-; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-NEXT: .reg .b32 %f<7>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
@@ -516,7 +516,7 @@
; SM90: {
; SM90-NEXT: .reg .b16 %rs<5>;
; SM90-NEXT: .reg .b32 %r<4>;
-; SM90-NEXT: .reg .f32 %f<7>;
+; SM90-NEXT: .reg .b32 %f<7>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
; SM90-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
@@ -566,7 +566,7 @@
; SM70-LABEL: test_fpext_float(
; SM70: {
; SM70-NEXT: .reg .b32 %r<3>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fpext_float_param_0];
@@ -578,7 +578,7 @@
; SM80-LABEL: test_fpext_float(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];
@@ -589,7 +589,7 @@
; SM80-FTZ-LABEL: test_fpext_float(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];
@@ -600,7 +600,7 @@
; SM90-LABEL: test_fpext_float(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
-; SM90-NEXT: .reg .f32 %f<2>;
+; SM90-NEXT: .reg .b32 %f<2>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
; SM90-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];
@@ -617,7 +617,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.f32 %f1, [test_fptrunc_float_param_0];
@@ -635,7 +635,7 @@
; SM80-LABEL: test_fptrunc_float(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.f32 %f1, [test_fptrunc_float_param_0];
@@ -646,7 +646,7 @@
; SM80-FTZ-LABEL: test_fptrunc_float(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.f32 %f1, [test_fptrunc_float_param_0];
@@ -657,7 +657,7 @@
; SM90-LABEL: test_fptrunc_float(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
-; SM90-NEXT: .reg .f32 %f<2>;
+; SM90-NEXT: .reg .b32 %f<2>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
; SM90-NEXT: ld.param.f32 %f1, [test_fptrunc_float_param_0];
@@ -674,7 +674,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<9>;
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fadd_imm_1_param_0];
@@ -706,7 +706,7 @@
; SM80-FTZ-LABEL: test_fadd_imm_1(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
-; SM80-FTZ-NEXT: .reg .f32 %f<3>;
+; SM80-FTZ-NEXT: .reg .b32 %f<3>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];
@@ -735,7 +735,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [test_select_cc_bf16_f64_param_0];
@@ -756,7 +756,7 @@
; SM70: {
; SM70-NEXT: .reg .b16 %rs<9>;
; SM70-NEXT: .reg .b32 %r<21>;
-; SM70-NEXT: .reg .f32 %f<9>;
+; SM70-NEXT: .reg .b32 %f<9>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -798,7 +798,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<9>;
; SM80-NEXT: .reg .b32 %r<5>;
-; SM80-NEXT: .reg .f32 %f<9>;
+; SM80-NEXT: .reg .b32 %f<9>;
; SM80-NEXT: .reg .b64 %rd<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
@@ -824,7 +824,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<9>;
; SM80-FTZ-NEXT: .reg .b32 %r<5>;
-; SM80-FTZ-NEXT: .reg .f32 %f<9>;
+; SM80-FTZ-NEXT: .reg .b32 %f<9>;
; SM80-FTZ-NEXT: .reg .b64 %rd<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
@@ -850,7 +850,7 @@
; SM90: {
; SM90-NEXT: .reg .b16 %rs<9>;
; SM90-NEXT: .reg .b32 %r<5>;
-; SM90-NEXT: .reg .f32 %f<9>;
+; SM90-NEXT: .reg .b32 %f<9>;
; SM90-NEXT: .reg .b64 %rd<2>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
@@ -881,7 +881,7 @@
; SM70: {
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<4>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fptosi_i16_param_0];
@@ -896,7 +896,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];
@@ -910,7 +910,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
; SM80-FTZ-NEXT: .reg .b32 %r<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];
@@ -940,7 +940,7 @@
; SM70: {
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<4>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_fptoui_i16_param_0];
@@ -955,7 +955,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];
@@ -969,7 +969,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
; SM80-FTZ-NEXT: .reg .b32 %r<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];
@@ -1000,7 +1000,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<3>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %rs1, [test_sitofp_i16_param_0];
@@ -1019,7 +1019,7 @@
; SM80-LABEL: test_sitofp_i16(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.u16 %rs1, [test_sitofp_i16_param_0];
@@ -1031,7 +1031,7 @@
; SM80-FTZ-LABEL: test_sitofp_i16(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.u16 %rs1, [test_sitofp_i16_param_0];
@@ -1059,7 +1059,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<3>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u8 %rs1, [test_uitofp_i8_param_0];
@@ -1078,7 +1078,7 @@
; SM80-LABEL: test_uitofp_i8(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.u8 %rs1, [test_uitofp_i8_param_0];
@@ -1090,7 +1090,7 @@
; SM80-FTZ-LABEL: test_uitofp_i8(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.u8 %rs1, [test_uitofp_i8_param_0];
@@ -1118,7 +1118,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<4>;
; SM70-NEXT: .reg .b32 %r<8>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u8 %rs1, [test_uitofp_i1_param_0];
@@ -1142,7 +1142,7 @@
; SM80-NEXT: .reg .pred %p<2>;
; SM80-NEXT: .reg .b16 %rs<4>;
; SM80-NEXT: .reg .b32 %r<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.u8 %rs1, [test_uitofp_i1_param_0];
@@ -1159,7 +1159,7 @@
; SM80-FTZ-NEXT: .reg .pred %p<2>;
; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
; SM80-FTZ-NEXT: .reg .b32 %r<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.u8 %rs1, [test_uitofp_i1_param_0];
@@ -1195,7 +1195,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<3>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %rs1, [test_uitofp_i16_param_0];
@@ -1214,7 +1214,7 @@
; SM80-LABEL: test_uitofp_i16(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.u16 %rs1, [test_uitofp_i16_param_0];
@@ -1226,7 +1226,7 @@
; SM80-FTZ-LABEL: test_uitofp_i16(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.u16 %rs1, [test_uitofp_i16_param_0];
@@ -1254,7 +1254,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<8>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u32 %r1, [test_uitofp_i32_param_0];
@@ -1274,7 +1274,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<2>;
; SM80-NEXT: .reg .b32 %r<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.u32 %r1, [test_uitofp_i32_param_0];
@@ -1287,7 +1287,7 @@
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
; SM80-FTZ-NEXT: .reg .b32 %r<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.u32 %r1, [test_uitofp_i32_param_0];
@@ -1316,7 +1316,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<2>;
+; SM70-NEXT: .reg .b32 %f<2>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -1336,7 +1336,7 @@
; SM80-LABEL: test_uitofp_i64(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<2>;
-; SM80-NEXT: .reg .f32 %f<2>;
+; SM80-NEXT: .reg .b32 %f<2>;
; SM80-NEXT: .reg .b64 %rd<2>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
@@ -1349,7 +1349,7 @@
; SM80-FTZ-LABEL: test_uitofp_i64(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
-; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b32 %f<2>;
; SM80-FTZ-NEXT: .reg .b64 %rd<2>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
@@ -1379,7 +1379,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<9>;
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_roundeven_param_0];
@@ -1400,7 +1400,7 @@
; SM80-LABEL: test_roundeven(
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
-; SM80-NEXT: .reg .f32 %f<3>;
+; SM80-NEXT: .reg .b32 %f<3>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];
@@ -1413,7 +1413,7 @@
; SM80-FTZ-LABEL: test_roundeven(
; SM80-FTZ: {
; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
-; SM80-FTZ-NEXT: .reg .f32 %f<3>;
+; SM80-FTZ-NEXT: .reg .b32 %f<3>;
; SM80-FTZ-EMPTY:
; SM80-FTZ-NEXT: // %bb.0:
; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];
@@ -1442,7 +1442,7 @@
; SM70-NEXT: .reg .pred %p<6>;
; SM70-NEXT: .reg .b16 %rs<8>;
; SM70-NEXT: .reg .b32 %r<7>;
-; SM70-NEXT: .reg .f32 %f<4>;
+; SM70-NEXT: .reg .b32 %f<4>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];
@@ -1511,7 +1511,7 @@
; SM70-NEXT: .reg .pred %p<2>;
; SM70-NEXT: .reg .b16 %rs<2>;
; SM70-NEXT: .reg .b32 %r<11>;
-; SM70-NEXT: .reg .f32 %f<4>;
+; SM70-NEXT: .reg .b32 %f<4>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u16 %r1, [test_maxnum_param_1];
@@ -1574,7 +1574,7 @@
; SM70-NEXT: .reg .pred %p<11>;
; SM70-NEXT: .reg .b16 %rs<15>;
; SM70-NEXT: .reg .b32 %r<16>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_0];
@@ -1665,7 +1665,7 @@
; SM70-NEXT: .reg .pred %p<3>;
; SM70-NEXT: .reg .b16 %rs<5>;
; SM70-NEXT: .reg .b32 %r<24>;
-; SM70-NEXT: .reg .f32 %f<7>;
+; SM70-NEXT: .reg .b32 %f<7>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_0];
@@ -1741,4 +1741,4 @@
}
declare bfloat @llvm.maximum.bf16(bfloat, bfloat)
-declare <2 x bfloat> @llvm.maximum.v2bf16(<2 x bfloat>, <2 x bfloat>)
\ No newline at end of file
+declare <2 x bfloat> @llvm.maximum.v2bf16(<2 x bfloat>, <2 x bfloat>)
diff --git a/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll b/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll
index fdf481e..5ab684a 100644
--- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll
+++ b/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll
@@ -12,7 +12,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sin_param_0];
@@ -33,7 +33,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_cos_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
index 706a88f..677f0d7 100644
--- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
@@ -131,7 +131,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<7>;
+; CHECK-NEXT: .reg .b32 %f<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
@@ -259,7 +259,7 @@
; SM80-NEXT: .reg .pred %p<3>;
; SM80-NEXT: .reg .b16 %rs<11>;
; SM80-NEXT: .reg .b32 %r<6>;
-; SM80-NEXT: .reg .f32 %f<5>;
+; SM80-NEXT: .reg .b32 %f<5>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
@@ -312,7 +312,7 @@
; SM80-NEXT: .reg .pred %p<3>;
; SM80-NEXT: .reg .b16 %rs<5>;
; SM80-NEXT: .reg .b32 %r<3>;
-; SM80-NEXT: .reg .f32 %f<11>;
+; SM80-NEXT: .reg .b32 %f<11>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_select_cc_f32_bf16_param_0];
@@ -336,7 +336,7 @@
; SM90: {
; SM90-NEXT: .reg .pred %p<3>;
; SM90-NEXT: .reg .b32 %r<3>;
-; SM90-NEXT: .reg .f32 %f<7>;
+; SM90-NEXT: .reg .b32 %f<7>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
; SM90-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_select_cc_f32_bf16_param_0];
@@ -360,7 +360,7 @@
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_bf16_f32_param_0];
@@ -386,7 +386,7 @@
; CHECK-LABEL: test_fptrunc_2xfloat(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_fptrunc_2xfloat_param_0];
@@ -402,7 +402,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xfloat_param_0];
@@ -469,7 +469,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_param_0];
@@ -583,7 +583,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<3>;
-; SM80-NEXT: .reg .f32 %f<5>;
+; SM80-NEXT: .reg .b32 %f<5>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_floor_param_0];
@@ -618,7 +618,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<3>;
-; SM80-NEXT: .reg .f32 %f<5>;
+; SM80-NEXT: .reg .b32 %f<5>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_ceil_param_0];
@@ -653,7 +653,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<3>;
-; SM80-NEXT: .reg .f32 %f<5>;
+; SM80-NEXT: .reg .b32 %f<5>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_trunc_param_0];
@@ -688,7 +688,7 @@
; SM80: {
; SM80-NEXT: .reg .b16 %rs<3>;
; SM80-NEXT: .reg .b32 %r<3>;
-; SM80-NEXT: .reg .f32 %f<5>;
+; SM80-NEXT: .reg .b32 %f<5>;
; SM80-EMPTY:
; SM80-NEXT: // %bb.0:
; SM80-NEXT: ld.param.b32 %r1, [test_rint_param_0];
@@ -724,7 +724,7 @@
; CHECK-NEXT: .reg .pred %p<5>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<9>;
-; CHECK-NEXT: .reg .f32 %f<17>;
+; CHECK-NEXT: .reg .b32 %f<17>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_round_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll b/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll
index 93da3913..670e112 100644
--- a/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll
@@ -8,7 +8,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_u8_f32_param_0];
@@ -25,7 +25,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [cvt_u8_f64_param_0];
@@ -41,7 +41,7 @@
; CHECK-LABEL: cvt_f32_i8(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u8 %rs1, [cvt_f32_i8_param_0];
@@ -56,7 +56,7 @@
; CHECK-LABEL: cvt_f64_i8(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u8 %rs1, [cvt_f64_i8_param_0];
@@ -71,7 +71,7 @@
; CHECK-LABEL: cvt_f32_s8(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.s8 %rs1, [cvt_f32_s8_param_0];
@@ -86,7 +86,7 @@
; CHECK-LABEL: cvt_f64_s8(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.s8 %rs1, [cvt_f64_s8_param_0];
@@ -102,7 +102,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_s8_f32_param_0];
@@ -120,7 +120,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [cvt_s8_f64_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm100.ll b/llvm/test/CodeGen/NVPTX/convert-sm100.ll
index f92822f..7230872 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm100.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm100.ll
@@ -11,7 +11,7 @@
; CHECK-LABEL: cvt_rn_satf_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_satf_tf32_f32_param_0];
@@ -26,7 +26,7 @@
; CHECK-LABEL: cvt_rn_relu_satf_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_satf_tf32_f32_param_0];
@@ -41,7 +41,7 @@
; CHECK-LABEL: cvt_rz_satf_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_satf_tf32_f32_param_0];
@@ -56,7 +56,7 @@
; CHECK-LABEL: cvt_rz_relu_satf_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_relu_satf_tf32_f32_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm100a.ll b/llvm/test/CodeGen/NVPTX/convert-sm100a.ll
index f0dd5f0..04d7a65 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm100a.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm100a.ll
@@ -11,7 +11,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_sf_e2m3x2_f32_param_0];
@@ -29,7 +29,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_sf_e2m3x2_f32_param_0];
@@ -47,7 +47,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_sf_e3m2x2_f32_param_0];
@@ -65,7 +65,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_sf_e3m2x2_f32_param_0];
@@ -143,7 +143,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_ue8m0x2_f32_param_0];
@@ -161,7 +161,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_sf_ue8m0x2_f32_param_0];
@@ -179,7 +179,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rp_ue8m0x2_f32_param_0];
@@ -197,7 +197,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rp_sf_ue8m0x2_f32_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm80.ll b/llvm/test/CodeGen/NVPTX/convert-sm80.ll
index aebc28b..eb7a6bd 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm80.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm80.ll
@@ -7,7 +7,7 @@
; CHECK-LABEL: cvt_rn_bf16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_bf16x2_f32_param_0];
@@ -23,7 +23,7 @@
; CHECK-LABEL: cvt_rn_relu_bf16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_bf16x2_f32_param_0];
@@ -39,7 +39,7 @@
; CHECK-LABEL: cvt_rz_bf16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_bf16x2_f32_param_0];
@@ -55,7 +55,7 @@
; CHECK-LABEL: cvt_rz_relu_bf16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_relu_bf16x2_f32_param_0];
@@ -76,7 +76,7 @@
; CHECK-LABEL: cvt_rn_f16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_f16x2_f32_param_0];
@@ -92,7 +92,7 @@
; CHECK-LABEL: cvt_rn_relu_f16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_f16x2_f32_param_0];
@@ -108,7 +108,7 @@
; CHECK-LABEL: cvt_rz_f16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_f16x2_f32_param_0];
@@ -124,7 +124,7 @@
; CHECK-LABEL: cvt_rz_relu_f16x2_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_relu_f16x2_f32_param_0];
@@ -145,7 +145,7 @@
; CHECK-LABEL: cvt_rn_bf16_f32(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_bf16_f32_param_0];
@@ -160,7 +160,7 @@
; CHECK-LABEL: cvt_rn_relu_bf16_f32(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_bf16_f32_param_0];
@@ -175,7 +175,7 @@
; CHECK-LABEL: cvt_rz_bf16_f32(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_bf16_f32_param_0];
@@ -190,7 +190,7 @@
; CHECK-LABEL: cvt_rz_relu_bf16_f32(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_relu_bf16_f32_param_0];
@@ -210,7 +210,7 @@
; CHECK-LABEL: cvt_rna_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rna_tf32_f32_param_0];
@@ -228,7 +228,7 @@
; CHECK-LABEL: fold_ff2bf16x2(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fold_ff2bf16x2_param_0];
@@ -247,7 +247,7 @@
; CHECK-LABEL: fold_ff2f16x2(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fold_ff2f16x2_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm90.ll b/llvm/test/CodeGen/NVPTX/convert-sm90.ll
index 5f610e0..340117f 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm90.ll
@@ -11,7 +11,7 @@
; CHECK-LABEL: cvt_rn_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_tf32_f32_param_0];
@@ -26,7 +26,7 @@
; CHECK-LABEL: cvt_rn_relu_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rn_relu_tf32_f32_param_0];
@@ -41,7 +41,7 @@
; CHECK-LABEL: cvt_rz_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_tf32_f32_param_0];
@@ -56,7 +56,7 @@
; CHECK-LABEL: cvt_rz_relu_tf32_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [cvt_rz_relu_tf32_f32_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/copysign.ll b/llvm/test/CodeGen/NVPTX/copysign.ll
index 843ef4d..2e305e68 100644
--- a/llvm/test/CodeGen/NVPTX/copysign.ll
+++ b/llvm/test/CodeGen/NVPTX/copysign.ll
@@ -8,7 +8,7 @@
define float @fcopysign_f_f(float %a, float %b) {
; CHECK-LABEL: fcopysign_f_f(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_f_param_0];
@@ -23,7 +23,7 @@
define double @fcopysign_d_d(double %a, double %b) {
; CHECK-LABEL: fcopysign_d_d(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_d_param_0];
@@ -39,7 +39,7 @@
; CHECK-LABEL: fcopysign_f_d(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -63,7 +63,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_h_param_0];
@@ -86,7 +86,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_f_param_0];
@@ -109,7 +109,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_h_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll b/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll
index 92b7379..a233616 100644
--- a/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll
+++ b/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll
@@ -59,9 +59,9 @@
; CHECK-LABEL: test_distributed_shared_cluster_float_atomic(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.u64 %rd1, [test_distributed_shared_cluster_float_atomic_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/div.ll b/llvm/test/CodeGen/NVPTX/div.ll
index 4f9d587..f8711e3 100644
--- a/llvm/test/CodeGen/NVPTX/div.ll
+++ b/llvm/test/CodeGen/NVPTX/div.ll
@@ -5,7 +5,7 @@
define float @div_full(float %a, float %b) {
; CHECK-LABEL: div_full(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [div_full_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/f16-abs.ll b/llvm/test/CodeGen/NVPTX/f16-abs.ll
index d12653e..d3aaedf 100644
--- a/llvm/test/CodeGen/NVPTX/f16-abs.ll
+++ b/llvm/test/CodeGen/NVPTX/f16-abs.ll
@@ -49,7 +49,7 @@
; CHECK-NOF16-LABEL: test_fabs(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [test_fabs_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
index e9edabd..e854e5a 100644
--- a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
@@ -109,7 +109,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fadd_param_1];
@@ -148,7 +148,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];
@@ -182,7 +182,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_1_param_0];
@@ -216,7 +216,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fsub_param_1];
@@ -254,7 +254,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<6>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<6>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fneg_param_0];
@@ -289,7 +289,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fmul_param_1];
@@ -316,7 +316,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<7>;
+; CHECK-NEXT: .reg .b32 %f<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r2, [test_fdiv_param_1];
@@ -351,7 +351,7 @@
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<15>;
+; CHECK-NEXT: .reg .b32 %f<15>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r2, [test_frem_param_1];
@@ -591,7 +591,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<11>;
; CHECK-NOF16-NEXT: .reg .b32 %r<6>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
@@ -623,7 +623,7 @@
; CHECK-F16: {
; CHECK-F16-NEXT: .reg .pred %p<3>;
; CHECK-F16-NEXT: .reg .b32 %r<3>;
-; CHECK-F16-NEXT: .reg .f32 %f<7>;
+; CHECK-F16-NEXT: .reg .b32 %f<7>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.v2.f32 {%f3, %f4}, [test_select_cc_f32_f16_param_1];
@@ -641,7 +641,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<11>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<11>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.f32 {%f3, %f4}, [test_select_cc_f32_f16_param_1];
@@ -672,7 +672,7 @@
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.f32 {%f3, %f4}, [test_select_cc_f16_f32_param_3];
@@ -716,7 +716,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_une_param_1];
@@ -760,7 +760,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ueq_param_1];
@@ -804,7 +804,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ugt_param_1];
@@ -848,7 +848,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uge_param_1];
@@ -892,7 +892,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ult_param_1];
@@ -936,7 +936,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ule_param_1];
@@ -981,7 +981,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uno_param_1];
@@ -1025,7 +1025,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_one_param_1];
@@ -1069,7 +1069,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oeq_param_1];
@@ -1113,7 +1113,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ogt_param_1];
@@ -1157,7 +1157,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oge_param_1];
@@ -1201,7 +1201,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_olt_param_1];
@@ -1245,7 +1245,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ole_param_1];
@@ -1289,7 +1289,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ord_param_1];
@@ -1472,7 +1472,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];
@@ -1516,7 +1516,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_sitofp_2xi32_fadd_param_0];
@@ -1545,7 +1545,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_fptrunc_2xfloat_param_0];
@@ -1563,7 +1563,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.f64 {%fd1, %fd2}, [test_fptrunc_2xdouble_param_0];
@@ -1581,7 +1581,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xfloat_param_0];
@@ -1599,7 +1599,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xdouble_param_0];
@@ -1643,7 +1643,7 @@
; CHECK-LABEL: test_bitcast_float_to_2xhalf(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [test_bitcast_float_to_2xhalf_param_0];
@@ -1658,7 +1658,7 @@
; CHECK-LABEL: test_bitcast_2xhalf_to_float(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_bitcast_2xhalf_to_float_param_0];
@@ -1698,7 +1698,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_param_0];
@@ -1728,7 +1728,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sin_param_0];
@@ -1751,7 +1751,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_cos_param_0];
@@ -1829,7 +1829,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;
; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<9>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<9>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fma_param_2];
@@ -1870,7 +1870,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fabs_param_0];
@@ -1893,7 +1893,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<7>;
+; CHECK-NEXT: .reg .b32 %f<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r2, [test_minnum_param_1];
@@ -1920,7 +1920,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<7>;
+; CHECK-NEXT: .reg .b32 %f<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r2, [test_maxnum_param_1];
@@ -1984,7 +1984,7 @@
; CHECK-F16: {
; CHECK-F16-NEXT: .reg .b16 %rs<3>;
; CHECK-F16-NEXT: .reg .b32 %r<6>;
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_copysign_f32_param_1];
@@ -2002,7 +2002,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.f32 {%f1, %f2}, [test_copysign_f32_param_1];
@@ -2031,7 +2031,7 @@
; CHECK-F16: {
; CHECK-F16-NEXT: .reg .b16 %rs<3>;
; CHECK-F16-NEXT: .reg .b32 %r<6>;
-; CHECK-F16-NEXT: .reg .f64 %fd<3>;
+; CHECK-F16-NEXT: .reg .b64 %fd<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.v2.f64 {%fd1, %fd2}, [test_copysign_f64_param_1];
@@ -2050,7 +2050,7 @@
; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
; CHECK-NOF16-NEXT: .reg .b64 %rd<7>;
-; CHECK-NOF16-NEXT: .reg .f64 %fd<3>;
+; CHECK-NOF16-NEXT: .reg .b64 %fd<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.f64 {%fd1, %fd2}, [test_copysign_f64_param_1];
@@ -2081,7 +2081,7 @@
; CHECK-F16: {
; CHECK-F16-NEXT: .reg .b16 %rs<3>;
; CHECK-F16-NEXT: .reg .b32 %r<6>;
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.b32 %r2, [test_copysign_extended_param_1];
@@ -2099,7 +2099,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<11>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_copysign_extended_param_1];
@@ -2236,7 +2236,7 @@
; CHECK-NEXT: .reg .pred %p<5>;
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<9>;
-; CHECK-NEXT: .reg .f32 %f<17>;
+; CHECK-NEXT: .reg .b32 %f<17>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_round_param_0];
@@ -2293,7 +2293,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;
; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<9>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<9>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fmuladd_param_2];
diff --git a/llvm/test/CodeGen/NVPTX/f32-ex2.ll b/llvm/test/CodeGen/NVPTX/f32-ex2.ll
index c9eff2a..2c5c814 100644
--- a/llvm/test/CodeGen/NVPTX/f32-ex2.ll
+++ b/llvm/test/CodeGen/NVPTX/f32-ex2.ll
@@ -9,7 +9,7 @@
define float @ex2_float(float %0) {
; CHECK-LABEL: ex2_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [ex2_float_param_0];
@@ -24,7 +24,7 @@
define float @ex2_float_ftz(float %0) {
; CHECK-LABEL: ex2_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [ex2_float_ftz_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/f32-lg2.ll b/llvm/test/CodeGen/NVPTX/f32-lg2.ll
index 43c5219..9dac308 100644
--- a/llvm/test/CodeGen/NVPTX/f32-lg2.ll
+++ b/llvm/test/CodeGen/NVPTX/f32-lg2.ll
@@ -10,7 +10,7 @@
define float @lg2_float(float %0) {
; CHECK-LABEL: lg2_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [lg2_float_param_0];
@@ -25,7 +25,7 @@
define float @lg2_float_ftz(float %0) {
; CHECK-LABEL: lg2_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [lg2_float_ftz_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll b/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll
index dd9ef22..d9c5a52 100644
--- a/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll
@@ -18,7 +18,7 @@
define float @fabs_float(float %a) {
; CHECK-LABEL: fabs_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fabs_float_param_0];
@@ -32,7 +32,7 @@
define float @fabs_float_ftz(float %a) {
; CHECK-LABEL: fabs_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fabs_float_ftz_param_0];
@@ -46,7 +46,7 @@
define double @fabs_double(double %a) {
; CHECK-LABEL: fabs_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [fabs_double_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fexp2.ll b/llvm/test/CodeGen/NVPTX/fexp2.ll
index 7e485dc..4664d70 100644
--- a/llvm/test/CodeGen/NVPTX/fexp2.ll
+++ b/llvm/test/CodeGen/NVPTX/fexp2.ll
@@ -13,7 +13,7 @@
define float @exp2_test(float %in) {
; CHECK-LABEL: exp2_test(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.f32 %f1, [exp2_test_param_0];
@@ -23,7 +23,7 @@
;
; CHECK-FP16-LABEL: exp2_test(
; CHECK-FP16: {
-; CHECK-FP16-NEXT: .reg .f32 %f<3>;
+; CHECK-FP16-NEXT: .reg .b32 %f<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.f32 %f1, [exp2_test_param_0];
@@ -33,7 +33,7 @@
;
; CHECK-BF16-LABEL: exp2_test(
; CHECK-BF16: {
-; CHECK-BF16-NEXT: .reg .f32 %f<3>;
+; CHECK-BF16-NEXT: .reg .b32 %f<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.f32 %f1, [exp2_test_param_0];
@@ -49,7 +49,7 @@
define float @exp2_ftz_test(float %in) #0 {
; CHECK-LABEL: exp2_ftz_test(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.f32 %f1, [exp2_ftz_test_param_0];
@@ -59,7 +59,7 @@
;
; CHECK-FP16-LABEL: exp2_ftz_test(
; CHECK-FP16: {
-; CHECK-FP16-NEXT: .reg .f32 %f<3>;
+; CHECK-FP16-NEXT: .reg .b32 %f<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.f32 %f1, [exp2_ftz_test_param_0];
@@ -69,7 +69,7 @@
;
; CHECK-BF16-LABEL: exp2_ftz_test(
; CHECK-BF16: {
-; CHECK-BF16-NEXT: .reg .f32 %f<3>;
+; CHECK-BF16-NEXT: .reg .b32 %f<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.f32 %f1, [exp2_ftz_test_param_0];
@@ -85,7 +85,7 @@
define <2 x float> @exp2_test_v(<2 x float> %in) {
; CHECK-LABEL: exp2_test_v(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.v2.f32 {%f1, %f2}, [exp2_test_v_param_0];
@@ -96,7 +96,7 @@
;
; CHECK-FP16-LABEL: exp2_test_v(
; CHECK-FP16: {
-; CHECK-FP16-NEXT: .reg .f32 %f<5>;
+; CHECK-FP16-NEXT: .reg .b32 %f<5>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.v2.f32 {%f1, %f2}, [exp2_test_v_param_0];
@@ -107,7 +107,7 @@
;
; CHECK-BF16-LABEL: exp2_test_v(
; CHECK-BF16: {
-; CHECK-BF16-NEXT: .reg .f32 %f<5>;
+; CHECK-BF16-NEXT: .reg .b32 %f<5>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.v2.f32 {%f1, %f2}, [exp2_test_v_param_0];
@@ -127,7 +127,7 @@
; CHECK-LABEL: exp2_f16_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];
@@ -167,7 +167,7 @@
; CHECK-LABEL: exp2_f16_ftz_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];
@@ -207,7 +207,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [exp2_f16_test_v_param_0];
@@ -256,7 +256,7 @@
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<9>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.u16 %r1, [exp2_bf16_test_param_0];
@@ -279,7 +279,7 @@
; CHECK-FP16-NEXT: .reg .pred %p<2>;
; CHECK-FP16-NEXT: .reg .b16 %rs<2>;
; CHECK-FP16-NEXT: .reg .b32 %r<9>;
-; CHECK-FP16-NEXT: .reg .f32 %f<3>;
+; CHECK-FP16-NEXT: .reg .b32 %f<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.u16 %r1, [exp2_bf16_test_param_0];
@@ -318,7 +318,7 @@
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<19>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [exp2_bf16_test_v_param_0];
@@ -354,7 +354,7 @@
; CHECK-FP16-NEXT: .reg .pred %p<3>;
; CHECK-FP16-NEXT: .reg .b16 %rs<3>;
; CHECK-FP16-NEXT: .reg .b32 %r<19>;
-; CHECK-FP16-NEXT: .reg .f32 %f<5>;
+; CHECK-FP16-NEXT: .reg .b32 %f<5>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_bf16_test_v_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/flog2.ll b/llvm/test/CodeGen/NVPTX/flog2.ll
index ff762dc..4dfed3d 100644
--- a/llvm/test/CodeGen/NVPTX/flog2.ll
+++ b/llvm/test/CodeGen/NVPTX/flog2.ll
@@ -7,7 +7,7 @@
define float @log2_test(float %in) {
; CHECK-LABEL: log2_test(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.f32 %f1, [log2_test_param_0];
@@ -23,7 +23,7 @@
define float @log2_ftz_test(float %in) #0 {
; CHECK-LABEL: log2_ftz_test(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.f32 %f1, [log2_ftz_test_param_0];
@@ -39,7 +39,7 @@
define <2 x float> @log2_test_v(<2 x float> %in) {
; CHECK-LABEL: log2_test_v(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.v2.f32 {%f1, %f2}, [log2_test_v_param_0];
@@ -59,7 +59,7 @@
; CHECK-LABEL: log2_f16_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [log2_f16_test_param_0];
@@ -78,7 +78,7 @@
; CHECK-LABEL: log2_f16_ftz_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [log2_f16_ftz_test_param_0];
@@ -98,7 +98,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [log2_f16_test_v_param_0];
@@ -126,7 +126,7 @@
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<9>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.u16 %r1, [log2_bf16_test_param_0];
@@ -155,7 +155,7 @@
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<9>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.u16 %r1, [log2_bf16_ftz_test_param_0];
@@ -184,7 +184,7 @@
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<19>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [log2_bf16_test_v_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll b/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll
index 7dce894..9051a0b 100644
--- a/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll
+++ b/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll
@@ -198,7 +198,7 @@
; CHECK-SM70-LABEL: fma_f16_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<3>;
+; CHECK-SM70-NEXT: .reg .b32 %f<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];
@@ -250,7 +250,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
-; CHECK-SM70-NEXT: .reg .f32 %f<6>;
+; CHECK-SM70-NEXT: .reg .b32 %f<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_unsafe_with_nans_param_2];
@@ -314,7 +314,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
-; CHECK-SM70-NEXT: .reg .f32 %f<6>;
+; CHECK-SM70-NEXT: .reg .b32 %f<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_no_nans_param_2];
@@ -372,7 +372,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<9>;
; CHECK-FTZ-NEXT: .reg .b32 %r<7>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<6>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<6>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];
@@ -402,7 +402,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<4>;
; CHECK-SM70-NEXT: .reg .b32 %r<29>;
-; CHECK-SM70-NEXT: .reg .f32 %f<10>;
+; CHECK-SM70-NEXT: .reg .b32 %f<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];
@@ -490,7 +490,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
-; CHECK-SM70-NEXT: .reg .f32 %f<7>;
+; CHECK-SM70-NEXT: .reg .b32 %f<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_maxnum_no_nans_param_2];
@@ -731,7 +731,7 @@
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<5>;
+; CHECK-SM70-NEXT: .reg .b32 %f<5>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_2];
@@ -788,7 +788,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<31>;
-; CHECK-SM70-NEXT: .reg .f32 %f<11>;
+; CHECK-SM70-NEXT: .reg .b32 %f<11>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_unsafe_with_nans_param_0];
@@ -881,7 +881,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<31>;
-; CHECK-SM70-NEXT: .reg .f32 %f<11>;
+; CHECK-SM70-NEXT: .reg .b32 %f<11>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];
@@ -968,7 +968,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<20>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<11>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<11>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
@@ -1012,7 +1012,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<9>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<61>;
-; CHECK-SM70-NEXT: .reg .f32 %f<19>;
+; CHECK-SM70-NEXT: .reg .b32 %f<19>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
@@ -1149,7 +1149,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<43>;
-; CHECK-SM70-NEXT: .reg .f32 %f<13>;
+; CHECK-SM70-NEXT: .reg .b32 %f<13>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fma-relu-fma-intrinsic.ll b/llvm/test/CodeGen/NVPTX/fma-relu-fma-intrinsic.ll
index eb51d7db..73f808f 100644
--- a/llvm/test/CodeGen/NVPTX/fma-relu-fma-intrinsic.ll
+++ b/llvm/test/CodeGen/NVPTX/fma-relu-fma-intrinsic.ll
@@ -137,7 +137,7 @@
; CHECK-SM70-LABEL: fma_f16_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<3>;
+; CHECK-SM70-NEXT: .reg .b32 %f<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];
@@ -184,7 +184,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
-; CHECK-SM70-NEXT: .reg .f32 %f<6>;
+; CHECK-SM70-NEXT: .reg .b32 %f<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_no_nans_param_2];
@@ -239,7 +239,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<5>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];
@@ -264,7 +264,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<4>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<27>;
-; CHECK-SM70-NEXT: .reg .f32 %f<9>;
+; CHECK-SM70-NEXT: .reg .b32 %f<9>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];
@@ -345,7 +345,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
-; CHECK-SM70-NEXT: .reg .f32 %f<7>;
+; CHECK-SM70-NEXT: .reg .b32 %f<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_maxnum_no_nans_param_2];
@@ -516,7 +516,7 @@
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<5>;
+; CHECK-SM70-NEXT: .reg .b32 %f<5>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_2];
@@ -568,7 +568,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<31>;
-; CHECK-SM70-NEXT: .reg .f32 %f<11>;
+; CHECK-SM70-NEXT: .reg .b32 %f<11>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];
@@ -652,7 +652,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-NEXT: .reg .b32 %r<14>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<9>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<9>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];
@@ -687,7 +687,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<7>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<57>;
-; CHECK-SM70-NEXT: .reg .f32 %f<17>;
+; CHECK-SM70-NEXT: .reg .b32 %f<17>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];
@@ -811,7 +811,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<43>;
-; CHECK-SM70-NEXT: .reg .f32 %f<13>;
+; CHECK-SM70-NEXT: .reg .b32 %f<13>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fma-relu-instruction-flag.ll b/llvm/test/CodeGen/NVPTX/fma-relu-instruction-flag.ll
index a3545f5..b94fa5a 100644
--- a/llvm/test/CodeGen/NVPTX/fma-relu-instruction-flag.ll
+++ b/llvm/test/CodeGen/NVPTX/fma-relu-instruction-flag.ll
@@ -147,7 +147,7 @@
; CHECK-SM70-LABEL: fma_f16_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<3>;
+; CHECK-SM70-NEXT: .reg .b32 %f<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];
@@ -195,7 +195,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
-; CHECK-SM70-NEXT: .reg .f32 %f<6>;
+; CHECK-SM70-NEXT: .reg .b32 %f<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_no_nans_param_2];
@@ -253,7 +253,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<9>;
; CHECK-FTZ-NEXT: .reg .b32 %r<7>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<6>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<6>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];
@@ -283,7 +283,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<4>;
; CHECK-SM70-NEXT: .reg .b32 %r<29>;
-; CHECK-SM70-NEXT: .reg .f32 %f<10>;
+; CHECK-SM70-NEXT: .reg .b32 %f<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];
@@ -373,7 +373,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
-; CHECK-SM70-NEXT: .reg .f32 %f<7>;
+; CHECK-SM70-NEXT: .reg .b32 %f<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_expanded_maxnum_no_nans_param_2];
@@ -563,7 +563,7 @@
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<5>;
+; CHECK-SM70-NEXT: .reg .b32 %f<5>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_2];
@@ -616,7 +616,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<31>;
-; CHECK-SM70-NEXT: .reg .f32 %f<11>;
+; CHECK-SM70-NEXT: .reg .b32 %f<11>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];
@@ -703,7 +703,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<20>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<11>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<11>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
@@ -747,7 +747,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<9>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<61>;
-; CHECK-SM70-NEXT: .reg .f32 %f<19>;
+; CHECK-SM70-NEXT: .reg .b32 %f<19>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
@@ -884,7 +884,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<43>;
-; CHECK-SM70-NEXT: .reg .f32 %f<13>;
+; CHECK-SM70-NEXT: .reg .b32 %f<13>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];
@@ -1084,7 +1084,7 @@
; CHECK-SM70-LABEL: fma_f16_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<3>;
+; CHECK-SM70-NEXT: .reg .b32 %f<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];
@@ -1131,7 +1131,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
-; CHECK-SM70-NEXT: .reg .f32 %f<6>;
+; CHECK-SM70-NEXT: .reg .b32 %f<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_no_nans_param_2];
@@ -1186,7 +1186,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<5>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];
@@ -1211,7 +1211,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<4>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<27>;
-; CHECK-SM70-NEXT: .reg .f32 %f<9>;
+; CHECK-SM70-NEXT: .reg .b32 %f<9>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];
@@ -1292,7 +1292,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
-; CHECK-SM70-NEXT: .reg .f32 %f<7>;
+; CHECK-SM70-NEXT: .reg .b32 %f<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.u16 %r1, [fma_bf16_maxnum_no_nans_param_2];
@@ -1467,7 +1467,7 @@
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
-; CHECK-SM70-NEXT: .reg .f32 %f<5>;
+; CHECK-SM70-NEXT: .reg .b32 %f<5>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_2];
@@ -1519,7 +1519,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<31>;
-; CHECK-SM70-NEXT: .reg .f32 %f<11>;
+; CHECK-SM70-NEXT: .reg .b32 %f<11>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];
@@ -1603,7 +1603,7 @@
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-NEXT: .reg .b32 %r<14>;
-; CHECK-FTZ-NEXT: .reg .f32 %f<9>;
+; CHECK-FTZ-NEXT: .reg .b32 %f<9>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];
@@ -1638,7 +1638,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<7>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<57>;
-; CHECK-SM70-NEXT: .reg .f32 %f<17>;
+; CHECK-SM70-NEXT: .reg .b32 %f<17>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];
@@ -1762,7 +1762,7 @@
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<43>;
-; CHECK-SM70-NEXT: .reg .f32 %f<13>;
+; CHECK-SM70-NEXT: .reg .b32 %f<13>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/fp-contract.ll b/llvm/test/CodeGen/NVPTX/fp-contract.ll
index ea5da6e..bd559ea 100644
--- a/llvm/test/CodeGen/NVPTX/fp-contract.ll
+++ b/llvm/test/CodeGen/NVPTX/fp-contract.ll
@@ -15,7 +15,7 @@
define float @t0(float %a, float %b, float %c) {
; FAST-LABEL: t0(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<5>;
+; FAST-NEXT: .reg .b32 %f<5>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [t0_param_0];
@@ -27,7 +27,7 @@
;
; DEFAULT-LABEL: t0(
; DEFAULT: {
-; DEFAULT-NEXT: .reg .f32 %f<6>;
+; DEFAULT-NEXT: .reg .b32 %f<6>;
; DEFAULT-EMPTY:
; DEFAULT-NEXT: // %bb.0:
; DEFAULT-NEXT: ld.param.f32 %f1, [t0_param_0];
@@ -47,7 +47,7 @@
define float @t1(float %a, float %b) {
; FAST-LABEL: t1(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<6>;
+; FAST-NEXT: .reg .b32 %f<6>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [t1_param_0];
@@ -60,7 +60,7 @@
;
; DEFAULT-LABEL: t1(
; DEFAULT: {
-; DEFAULT-NEXT: .reg .f32 %f<6>;
+; DEFAULT-NEXT: .reg .b32 %f<6>;
; DEFAULT-EMPTY:
; DEFAULT-NEXT: // %bb.0:
; DEFAULT-NEXT: ld.param.f32 %f1, [t1_param_0];
@@ -81,7 +81,7 @@
define float @t2(float %a, float %b) {
; CHECK-LABEL: t2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<6>;
+; CHECK-NEXT: .reg .b32 %f<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [t2_param_0];
@@ -101,7 +101,7 @@
define float @t3(float %a, float %b, float %c) {
; CHECK-LABEL: t3(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [t3_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/frem.ll b/llvm/test/CodeGen/NVPTX/frem.ll
index 73debfb..4077f6d 100644
--- a/llvm/test/CodeGen/NVPTX/frem.ll
+++ b/llvm/test/CodeGen/NVPTX/frem.ll
@@ -9,7 +9,7 @@
; FAST-LABEL: frem_f16(
; FAST: {
; FAST-NEXT: .reg .b16 %rs<4>;
-; FAST-NEXT: .reg .f32 %f<7>;
+; FAST-NEXT: .reg .b32 %f<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.b16 %rs1, [frem_f16_param_0];
@@ -28,7 +28,7 @@
; NORMAL: {
; NORMAL-NEXT: .reg .pred %p<2>;
; NORMAL-NEXT: .reg .b16 %rs<4>;
-; NORMAL-NEXT: .reg .f32 %f<8>;
+; NORMAL-NEXT: .reg .b32 %f<8>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.b16 %rs1, [frem_f16_param_0];
@@ -51,7 +51,7 @@
define float @frem_f32(float %a, float %b) {
; FAST-LABEL: frem_f32(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<7>;
+; FAST-NEXT: .reg .b32 %f<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_param_0];
@@ -66,7 +66,7 @@
; NORMAL-LABEL: frem_f32(
; NORMAL: {
; NORMAL-NEXT: .reg .pred %p<2>;
-; NORMAL-NEXT: .reg .f32 %f<8>;
+; NORMAL-NEXT: .reg .b32 %f<8>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_param_0];
@@ -86,7 +86,7 @@
define double @frem_f64(double %a, double %b) {
; FAST-LABEL: frem_f64(
; FAST: {
-; FAST-NEXT: .reg .f64 %fd<7>;
+; FAST-NEXT: .reg .b64 %fd<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f64 %fd1, [frem_f64_param_0];
@@ -101,7 +101,7 @@
; NORMAL-LABEL: frem_f64(
; NORMAL: {
; NORMAL-NEXT: .reg .pred %p<2>;
-; NORMAL-NEXT: .reg .f64 %fd<8>;
+; NORMAL-NEXT: .reg .b64 %fd<8>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f64 %fd1, [frem_f64_param_0];
@@ -122,7 +122,7 @@
; FAST-LABEL: frem_f16_ninf(
; FAST: {
; FAST-NEXT: .reg .b16 %rs<4>;
-; FAST-NEXT: .reg .f32 %f<7>;
+; FAST-NEXT: .reg .b32 %f<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_param_0];
@@ -140,7 +140,7 @@
; NORMAL-LABEL: frem_f16_ninf(
; NORMAL: {
; NORMAL-NEXT: .reg .b16 %rs<4>;
-; NORMAL-NEXT: .reg .f32 %f<7>;
+; NORMAL-NEXT: .reg .b32 %f<7>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_param_0];
@@ -161,7 +161,7 @@
define float @frem_f32_ninf(float %a, float %b) {
; FAST-LABEL: frem_f32_ninf(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<7>;
+; FAST-NEXT: .reg .b32 %f<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_ninf_param_0];
@@ -175,7 +175,7 @@
;
; NORMAL-LABEL: frem_f32_ninf(
; NORMAL: {
-; NORMAL-NEXT: .reg .f32 %f<7>;
+; NORMAL-NEXT: .reg .b32 %f<7>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_ninf_param_0];
@@ -193,7 +193,7 @@
define double @frem_f64_ninf(double %a, double %b) {
; FAST-LABEL: frem_f64_ninf(
; FAST: {
-; FAST-NEXT: .reg .f64 %fd<7>;
+; FAST-NEXT: .reg .b64 %fd<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f64 %fd1, [frem_f64_ninf_param_0];
@@ -207,7 +207,7 @@
;
; NORMAL-LABEL: frem_f64_ninf(
; NORMAL: {
-; NORMAL-NEXT: .reg .f64 %fd<7>;
+; NORMAL-NEXT: .reg .b64 %fd<7>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f64 %fd1, [frem_f64_ninf_param_0];
@@ -225,7 +225,7 @@
define float @frem_f32_imm1(float %a) {
; FAST-LABEL: frem_f32_imm1(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<5>;
+; FAST-NEXT: .reg .b32 %f<5>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_imm1_param_0];
@@ -237,7 +237,7 @@
;
; NORMAL-LABEL: frem_f32_imm1(
; NORMAL: {
-; NORMAL-NEXT: .reg .f32 %f<5>;
+; NORMAL-NEXT: .reg .b32 %f<5>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_imm1_param_0];
@@ -253,7 +253,7 @@
define float @frem_f32_imm2(float %a) {
; FAST-LABEL: frem_f32_imm2(
; FAST: {
-; FAST-NEXT: .reg .f32 %f<7>;
+; FAST-NEXT: .reg .b32 %f<7>;
; FAST-EMPTY:
; FAST-NEXT: // %bb.0:
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_imm2_param_0];
@@ -268,7 +268,7 @@
; NORMAL-LABEL: frem_f32_imm2(
; NORMAL: {
; NORMAL-NEXT: .reg .pred %p<2>;
-; NORMAL-NEXT: .reg .f32 %f<8>;
+; NORMAL-NEXT: .reg .b32 %f<8>;
; NORMAL-EMPTY:
; NORMAL-NEXT: // %bb.0:
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_imm2_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index 1b779db..65edcf2 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -1141,7 +1141,7 @@
; CHECK-LABEL: test_bitcast_float_to_4xi8(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [test_bitcast_float_to_4xi8_param_0];
@@ -1169,7 +1169,7 @@
; CHECK-LABEL: test_bitcast_4xi8_to_float(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_bitcast_4xi8_to_float_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/intrinsics.ll b/llvm/test/CodeGen/NVPTX/intrinsics.ll
index cc6af06..01c51bb 100644
--- a/llvm/test/CodeGen/NVPTX/intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/intrinsics.ll
@@ -7,7 +7,7 @@
define float @test_fabsf(float %f) {
; CHECK-LABEL: test_fabsf(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [test_fabsf_param_0];
@@ -21,7 +21,7 @@
define double @test_fabs(double %d) {
; CHECK-LABEL: test_fabs(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [test_fabs_param_0];
@@ -35,7 +35,7 @@
define float @test_nvvm_sqrt(float %a) {
; CHECK-LABEL: test_nvvm_sqrt(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [test_nvvm_sqrt_param_0];
@@ -49,7 +49,7 @@
define float @test_llvm_sqrt(float %a) {
; CHECK-LABEL: test_llvm_sqrt(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [test_llvm_sqrt_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/ldg-invariant.ll b/llvm/test/CodeGen/NVPTX/ldg-invariant.ll
index 16a0189..2fe2d28 100644
--- a/llvm/test/CodeGen/NVPTX/ldg-invariant.ll
+++ b/llvm/test/CodeGen/NVPTX/ldg-invariant.ll
@@ -27,7 +27,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<4>;
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -54,7 +54,7 @@
; CHECK-LABEL: ld_global_v4f16(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<8>;
-; CHECK-NEXT: .reg .f32 %f<10>;
+; CHECK-NEXT: .reg .b32 %f<10>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -92,7 +92,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<8>;
; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-NEXT: .reg .f32 %f<10>;
+; CHECK-NEXT: .reg .b32 %f<10>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
diff --git a/llvm/test/CodeGen/NVPTX/ldu-ldg.ll b/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
index 4c5c44a..2c1550a 100644
--- a/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
+++ b/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
@@ -104,7 +104,7 @@
define float @test_ldu_f32(ptr addrspace(1) %ptr) {
; CHECK-LABEL: test_ldu_f32(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -120,7 +120,7 @@
; CHECK-LABEL: test_ldu_f64(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f64_param_0];
@@ -241,7 +241,7 @@
define float @test_ldg_f32(ptr addrspace(1) %ptr) {
; CHECK-LABEL: test_ldg_f32(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -257,7 +257,7 @@
; CHECK-LABEL: test_ldg_f64(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f64_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/load-store-scalars.ll b/llvm/test/CodeGen/NVPTX/load-store-scalars.ll
index ed94cb4..cb2e247 100644
--- a/llvm/test/CodeGen/NVPTX/load-store-scalars.ll
+++ b/llvm/test/CodeGen/NVPTX/load-store-scalars.ll
@@ -91,7 +91,7 @@
define void @generic_float(ptr %a) {
; CHECK-LABEL: generic_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -110,7 +110,7 @@
; CHECK-LABEL: generic_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_double_param_0];
@@ -200,7 +200,7 @@
define void @generic_volatile_float(ptr %a) {
; CHECK-LABEL: generic_volatile_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -219,7 +219,7 @@
; CHECK-LABEL: generic_volatile_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_volatile_double_param_0];
@@ -356,7 +356,7 @@
define void @generic_unordered_sys_float(ptr %a) {
; SM60-LABEL: generic_unordered_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -368,7 +368,7 @@
;
; SM70-LABEL: generic_unordered_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -387,7 +387,7 @@
; SM60-LABEL: generic_unordered_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [generic_unordered_sys_double_param_0];
@@ -399,7 +399,7 @@
; SM70-LABEL: generic_unordered_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [generic_unordered_sys_double_param_0];
@@ -489,7 +489,7 @@
define void @generic_unordered_volatile_sys_float(ptr %a) {
; CHECK-LABEL: generic_unordered_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -508,7 +508,7 @@
; CHECK-LABEL: generic_unordered_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_unordered_volatile_sys_double_param_0];
@@ -645,7 +645,7 @@
define void @generic_monotonic_sys_float(ptr %a) {
; SM60-LABEL: generic_monotonic_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -657,7 +657,7 @@
;
; SM70-LABEL: generic_monotonic_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -676,7 +676,7 @@
; SM60-LABEL: generic_monotonic_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [generic_monotonic_sys_double_param_0];
@@ -688,7 +688,7 @@
; SM70-LABEL: generic_monotonic_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [generic_monotonic_sys_double_param_0];
@@ -778,7 +778,7 @@
define void @generic_monotonic_volatile_sys_float(ptr %a) {
; CHECK-LABEL: generic_monotonic_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -797,7 +797,7 @@
; CHECK-LABEL: generic_monotonic_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_monotonic_volatile_sys_double_param_0];
@@ -889,7 +889,7 @@
define void @global_float(ptr addrspace(1) %a) {
; CHECK-LABEL: global_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -908,7 +908,7 @@
; CHECK-LABEL: global_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [global_double_param_0];
@@ -998,7 +998,7 @@
define void @global_volatile_float(ptr addrspace(1) %a) {
; CHECK-LABEL: global_volatile_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1017,7 +1017,7 @@
; CHECK-LABEL: global_volatile_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [global_volatile_double_param_0];
@@ -1154,7 +1154,7 @@
define void @global_unordered_sys_float(ptr addrspace(1) %a) {
; SM60-LABEL: global_unordered_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -1166,7 +1166,7 @@
;
; SM70-LABEL: global_unordered_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -1185,7 +1185,7 @@
; SM60-LABEL: global_unordered_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [global_unordered_sys_double_param_0];
@@ -1197,7 +1197,7 @@
; SM70-LABEL: global_unordered_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [global_unordered_sys_double_param_0];
@@ -1334,7 +1334,7 @@
define void @global_unordered_volatile_sys_float(ptr addrspace(1) %a) {
; SM60-LABEL: global_unordered_volatile_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -1346,7 +1346,7 @@
;
; SM70-LABEL: global_unordered_volatile_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -1365,7 +1365,7 @@
; SM60-LABEL: global_unordered_volatile_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [global_unordered_volatile_sys_double_param_0];
@@ -1377,7 +1377,7 @@
; SM70-LABEL: global_unordered_volatile_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [global_unordered_volatile_sys_double_param_0];
@@ -1514,7 +1514,7 @@
define void @global_monotonic_sys_float(ptr addrspace(1) %a) {
; SM60-LABEL: global_monotonic_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -1526,7 +1526,7 @@
;
; SM70-LABEL: global_monotonic_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -1545,7 +1545,7 @@
; SM60-LABEL: global_monotonic_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [global_monotonic_sys_double_param_0];
@@ -1557,7 +1557,7 @@
; SM70-LABEL: global_monotonic_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [global_monotonic_sys_double_param_0];
@@ -1694,7 +1694,7 @@
define void @global_monotonic_volatile_sys_float(ptr addrspace(1) %a) {
; SM60-LABEL: global_monotonic_volatile_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -1706,7 +1706,7 @@
;
; SM70-LABEL: global_monotonic_volatile_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -1725,7 +1725,7 @@
; SM60-LABEL: global_monotonic_volatile_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [global_monotonic_volatile_sys_double_param_0];
@@ -1737,7 +1737,7 @@
; SM70-LABEL: global_monotonic_volatile_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [global_monotonic_volatile_sys_double_param_0];
@@ -1829,7 +1829,7 @@
define void @shared_float(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1848,7 +1848,7 @@
; CHECK-LABEL: shared_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_double_param_0];
@@ -1938,7 +1938,7 @@
define void @shared_volatile_float(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_volatile_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1957,7 +1957,7 @@
; CHECK-LABEL: shared_volatile_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_volatile_double_param_0];
@@ -2094,7 +2094,7 @@
define void @shared_unordered_sys_float(ptr addrspace(3) %a) {
; SM60-LABEL: shared_unordered_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -2106,7 +2106,7 @@
;
; SM70-LABEL: shared_unordered_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -2125,7 +2125,7 @@
; SM60-LABEL: shared_unordered_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [shared_unordered_sys_double_param_0];
@@ -2137,7 +2137,7 @@
; SM70-LABEL: shared_unordered_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [shared_unordered_sys_double_param_0];
@@ -2227,7 +2227,7 @@
define void @shared_unordered_volatile_sys_float(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_unordered_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2246,7 +2246,7 @@
; CHECK-LABEL: shared_unordered_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_unordered_volatile_sys_double_param_0];
@@ -2383,7 +2383,7 @@
define void @shared_monotonic_sys_float(ptr addrspace(3) %a) {
; SM60-LABEL: shared_monotonic_sys_float(
; SM60: {
-; SM60-NEXT: .reg .f32 %f<3>;
+; SM60-NEXT: .reg .b32 %f<3>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
@@ -2395,7 +2395,7 @@
;
; SM70-LABEL: shared_monotonic_sys_float(
; SM70: {
-; SM70-NEXT: .reg .f32 %f<3>;
+; SM70-NEXT: .reg .b32 %f<3>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
@@ -2414,7 +2414,7 @@
; SM60-LABEL: shared_monotonic_sys_double(
; SM60: {
; SM60-NEXT: .reg .b64 %rd<2>;
-; SM60-NEXT: .reg .f64 %fd<3>;
+; SM60-NEXT: .reg .b64 %fd<3>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
; SM60-NEXT: ld.param.u64 %rd1, [shared_monotonic_sys_double_param_0];
@@ -2426,7 +2426,7 @@
; SM70-LABEL: shared_monotonic_sys_double(
; SM70: {
; SM70-NEXT: .reg .b64 %rd<2>;
-; SM70-NEXT: .reg .f64 %fd<3>;
+; SM70-NEXT: .reg .b64 %fd<3>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
; SM70-NEXT: ld.param.u64 %rd1, [shared_monotonic_sys_double_param_0];
@@ -2516,7 +2516,7 @@
define void @shared_monotonic_volatile_sys_float(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_monotonic_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2535,7 +2535,7 @@
; CHECK-LABEL: shared_monotonic_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_monotonic_volatile_sys_double_param_0];
@@ -2627,7 +2627,7 @@
define void @local_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2646,7 +2646,7 @@
; CHECK-LABEL: local_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_double_param_0];
@@ -2736,7 +2736,7 @@
define void @local_volatile_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_volatile_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2755,7 +2755,7 @@
; CHECK-LABEL: local_volatile_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_volatile_double_param_0];
@@ -2845,7 +2845,7 @@
define void @local_unordered_sys_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_unordered_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2864,7 +2864,7 @@
; CHECK-LABEL: local_unordered_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_unordered_sys_double_param_0];
@@ -2954,7 +2954,7 @@
define void @local_unordered_volatile_sys_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_unordered_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2973,7 +2973,7 @@
; CHECK-LABEL: local_unordered_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_unordered_volatile_sys_double_param_0];
@@ -3063,7 +3063,7 @@
define void @local_monotonic_sys_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_monotonic_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -3082,7 +3082,7 @@
; CHECK-LABEL: local_monotonic_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_monotonic_sys_double_param_0];
@@ -3172,7 +3172,7 @@
define void @local_monotonic_volatile_sys_float(ptr addrspace(5) %a) {
; CHECK-LABEL: local_monotonic_volatile_sys_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -3191,7 +3191,7 @@
; CHECK-LABEL: local_monotonic_volatile_sys_double(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_monotonic_volatile_sys_double_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/load-store-vectors.ll b/llvm/test/CodeGen/NVPTX/load-store-vectors.ll
index ba397dc..3215fce 100644
--- a/llvm/test/CodeGen/NVPTX/load-store-vectors.ll
+++ b/llvm/test/CodeGen/NVPTX/load-store-vectors.ll
@@ -371,7 +371,7 @@
define void @generic_2xfloat(ptr %a) {
; CHECK-LABEL: generic_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -390,7 +390,7 @@
define void @generic_4xfloat(ptr %a) {
; CHECK-LABEL: generic_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -412,7 +412,7 @@
; CHECK-LABEL: generic_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_2xdouble_param_0];
@@ -792,7 +792,7 @@
define void @generic_volatile_2xfloat(ptr %a) {
; CHECK-LABEL: generic_volatile_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -811,7 +811,7 @@
define void @generic_volatile_4xfloat(ptr %a) {
; CHECK-LABEL: generic_volatile_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -833,7 +833,7 @@
; CHECK-LABEL: generic_volatile_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [generic_volatile_2xdouble_param_0];
@@ -1196,7 +1196,7 @@
define void @global_2xfloat(ptr addrspace(1) %a) {
; CHECK-LABEL: global_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1215,7 +1215,7 @@
define void @global_4xfloat(ptr addrspace(1) %a) {
; CHECK-LABEL: global_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1237,7 +1237,7 @@
; CHECK-LABEL: global_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [global_2xdouble_param_0];
@@ -1598,7 +1598,7 @@
define void @global_volatile_2xfloat(ptr addrspace(1) %a) {
; CHECK-LABEL: global_volatile_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1617,7 +1617,7 @@
define void @global_volatile_4xfloat(ptr addrspace(1) %a) {
; CHECK-LABEL: global_volatile_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -1639,7 +1639,7 @@
; CHECK-LABEL: global_volatile_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [global_volatile_2xdouble_param_0];
@@ -2002,7 +2002,7 @@
define void @shared_2xfloat(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2021,7 +2021,7 @@
define void @shared_4xfloat(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2043,7 +2043,7 @@
; CHECK-LABEL: shared_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_2xdouble_param_0];
@@ -2404,7 +2404,7 @@
define void @shared_volatile_2xfloat(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_volatile_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2423,7 +2423,7 @@
define void @shared_volatile_4xfloat(ptr addrspace(3) %a) {
; CHECK-LABEL: shared_volatile_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2445,7 +2445,7 @@
; CHECK-LABEL: shared_volatile_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [shared_volatile_2xdouble_param_0];
@@ -2808,7 +2808,7 @@
define void @local_2xfloat(ptr addrspace(5) %a) {
; CHECK-LABEL: local_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2827,7 +2827,7 @@
define void @local_4xfloat(ptr addrspace(5) %a) {
; CHECK-LABEL: local_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -2849,7 +2849,7 @@
; CHECK-LABEL: local_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_2xdouble_param_0];
@@ -3210,7 +3210,7 @@
define void @local_volatile_2xfloat(ptr addrspace(5) %a) {
; CHECK-LABEL: local_volatile_2xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -3229,7 +3229,7 @@
define void @local_volatile_4xfloat(ptr addrspace(5) %a) {
; CHECK-LABEL: local_volatile_4xfloat(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -3251,7 +3251,7 @@
; CHECK-LABEL: local_volatile_2xdouble(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [local_volatile_2xdouble_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/math-intrins.ll b/llvm/test/CodeGen/NVPTX/math-intrins.ll
index 189f342..a6d01c1 100644
--- a/llvm/test/CodeGen/NVPTX/math-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll
@@ -50,7 +50,7 @@
define float @ceil_float(float %a) {
; CHECK-LABEL: ceil_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [ceil_float_param_0];
@@ -64,7 +64,7 @@
define float @ceil_float_ftz(float %a) #1 {
; CHECK-LABEL: ceil_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [ceil_float_ftz_param_0];
@@ -78,7 +78,7 @@
define double @ceil_double(double %a) {
; CHECK-LABEL: ceil_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [ceil_double_param_0];
@@ -94,7 +94,7 @@
define float @floor_float(float %a) {
; CHECK-LABEL: floor_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [floor_float_param_0];
@@ -108,7 +108,7 @@
define float @floor_float_ftz(float %a) #1 {
; CHECK-LABEL: floor_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [floor_float_ftz_param_0];
@@ -122,7 +122,7 @@
define double @floor_double(double %a) {
; CHECK-LABEL: floor_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [floor_double_param_0];
@@ -141,7 +141,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [round_float_param_0];
@@ -169,7 +169,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-NEXT: .reg .b32 %f<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [round_float_ftz_param_0];
@@ -196,7 +196,7 @@
; CHECK-LABEL: round_double(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
-; CHECK-NEXT: .reg .f64 %fd<8>;
+; CHECK-NEXT: .reg .b64 %fd<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [round_double_param_0];
@@ -219,7 +219,7 @@
define float @nearbyint_float(float %a) {
; CHECK-LABEL: nearbyint_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [nearbyint_float_param_0];
@@ -233,7 +233,7 @@
define float @nearbyint_float_ftz(float %a) #1 {
; CHECK-LABEL: nearbyint_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [nearbyint_float_ftz_param_0];
@@ -247,7 +247,7 @@
define double @nearbyint_double(double %a) {
; CHECK-LABEL: nearbyint_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [nearbyint_double_param_0];
@@ -263,7 +263,7 @@
define float @rint_float(float %a) {
; CHECK-LABEL: rint_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [rint_float_param_0];
@@ -277,7 +277,7 @@
define float @rint_float_ftz(float %a) #1 {
; CHECK-LABEL: rint_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [rint_float_ftz_param_0];
@@ -291,7 +291,7 @@
define double @rint_double(double %a) {
; CHECK-LABEL: rint_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [rint_double_param_0];
@@ -307,7 +307,7 @@
define float @roundeven_float(float %a) {
; CHECK-LABEL: roundeven_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [roundeven_float_param_0];
@@ -321,7 +321,7 @@
define float @roundeven_float_ftz(float %a) #1 {
; CHECK-LABEL: roundeven_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [roundeven_float_ftz_param_0];
@@ -335,7 +335,7 @@
define double @roundeven_double(double %a) {
; CHECK-LABEL: roundeven_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [roundeven_double_param_0];
@@ -351,7 +351,7 @@
define float @trunc_float(float %a) {
; CHECK-LABEL: trunc_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [trunc_float_param_0];
@@ -365,7 +365,7 @@
define float @trunc_float_ftz(float %a) #1 {
; CHECK-LABEL: trunc_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [trunc_float_ftz_param_0];
@@ -379,7 +379,7 @@
define double @trunc_double(double %a) {
; CHECK-LABEL: trunc_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [trunc_double_param_0];
@@ -395,7 +395,7 @@
define float @abs_float(float %a) {
; CHECK-LABEL: abs_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [abs_float_param_0];
@@ -409,7 +409,7 @@
define float @abs_float_ftz(float %a) #1 {
; CHECK-LABEL: abs_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [abs_float_ftz_param_0];
@@ -423,7 +423,7 @@
define double @abs_double(double %a) {
; CHECK-LABEL: abs_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-NEXT: .reg .b64 %fd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [abs_double_param_0];
@@ -440,7 +440,7 @@
; CHECK-NOF16-LABEL: minnum_half(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];
@@ -466,7 +466,7 @@
; CHECK-SM80-NOF16-LABEL: minnum_half(
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];
@@ -484,7 +484,7 @@
define float @minnum_float(float %a, float %b) {
; CHECK-LABEL: minnum_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [minnum_float_param_0];
@@ -499,7 +499,7 @@
define float @minnum_imm1(float %a) {
; CHECK-LABEL: minnum_imm1(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [minnum_imm1_param_0];
@@ -513,7 +513,7 @@
define float @minnum_imm2(float %a) {
; CHECK-LABEL: minnum_imm2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [minnum_imm2_param_0];
@@ -527,7 +527,7 @@
define float @minnum_float_ftz(float %a, float %b) #1 {
; CHECK-LABEL: minnum_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [minnum_float_ftz_param_0];
@@ -542,7 +542,7 @@
define double @minnum_double(double %a, double %b) {
; CHECK-LABEL: minnum_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [minnum_double_param_0];
@@ -559,7 +559,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minnum_v2half_param_0];
@@ -593,7 +593,7 @@
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minnum_v2half_param_0];
@@ -622,7 +622,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<6>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<8>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];
@@ -658,7 +658,7 @@
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<8>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];
@@ -687,7 +687,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
@@ -708,7 +708,7 @@
;
; CHECK-F16-LABEL: minimum_float(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-NEXT: .reg .b32 %f<4>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
@@ -719,7 +719,7 @@
;
; CHECK-SM80-NOF16-LABEL: minimum_float(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
@@ -736,7 +736,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<4>;
; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<6>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<6>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
@@ -753,7 +753,7 @@
;
; CHECK-F16-LABEL: minimum_imm1(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
@@ -763,7 +763,7 @@
;
; CHECK-SM80-NOF16-LABEL: minimum_imm1(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
@@ -779,7 +779,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<4>;
; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<6>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<6>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
@@ -796,7 +796,7 @@
;
; CHECK-F16-LABEL: minimum_imm2(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
@@ -806,7 +806,7 @@
;
; CHECK-SM80-NOF16-LABEL: minimum_imm2(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
@@ -822,7 +822,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
@@ -843,7 +843,7 @@
;
; CHECK-F16-LABEL: minimum_float_ftz(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-NEXT: .reg .b32 %f<4>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
@@ -854,7 +854,7 @@
;
; CHECK-SM80-NOF16-LABEL: minimum_float_ftz(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
@@ -871,7 +871,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<5>;
; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-NEXT: .reg .f64 %fd<8>;
+; CHECK-NEXT: .reg .b64 %fd<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [minimum_double_param_0];
@@ -899,7 +899,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<11>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<15>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_v2half_param_0];
@@ -952,7 +952,7 @@
; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<15>;
; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_v2half_param_0];
@@ -998,7 +998,7 @@
; CHECK-NOF16-LABEL: maxnum_half(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];
@@ -1024,7 +1024,7 @@
; CHECK-SM80-NOF16-LABEL: maxnum_half(
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];
@@ -1042,7 +1042,7 @@
define float @maxnum_imm1(float %a) {
; CHECK-LABEL: maxnum_imm1(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [maxnum_imm1_param_0];
@@ -1056,7 +1056,7 @@
define float @maxnum_imm2(float %a) {
; CHECK-LABEL: maxnum_imm2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [maxnum_imm2_param_0];
@@ -1070,7 +1070,7 @@
define float @maxnum_float(float %a, float %b) {
; CHECK-LABEL: maxnum_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [maxnum_float_param_0];
@@ -1085,7 +1085,7 @@
define float @maxnum_float_ftz(float %a, float %b) #1 {
; CHECK-LABEL: maxnum_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [maxnum_float_ftz_param_0];
@@ -1100,7 +1100,7 @@
define double @maxnum_double(double %a, double %b) {
; CHECK-LABEL: maxnum_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [maxnum_double_param_0];
@@ -1117,7 +1117,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maxnum_v2half_param_0];
@@ -1151,7 +1151,7 @@
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;
; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maxnum_v2half_param_0];
@@ -1180,7 +1180,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<6>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<8>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];
@@ -1216,7 +1216,7 @@
; CHECK-SM80-NOF16: {
; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<8>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];
@@ -1244,7 +1244,7 @@
; CHECK-NOF16-LABEL: maximum_imm1(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
@@ -1258,7 +1258,7 @@
;
; CHECK-F16-LABEL: maximum_imm1(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
@@ -1268,7 +1268,7 @@
;
; CHECK-SM80-NOF16-LABEL: maximum_imm1(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
@@ -1283,7 +1283,7 @@
; CHECK-NOF16-LABEL: maximum_imm2(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
@@ -1297,7 +1297,7 @@
;
; CHECK-F16-LABEL: maximum_imm2(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-NEXT: .reg .b32 %f<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
@@ -1307,7 +1307,7 @@
;
; CHECK-SM80-NOF16-LABEL: maximum_imm2(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<3>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
@@ -1323,7 +1323,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
@@ -1344,7 +1344,7 @@
;
; CHECK-F16-LABEL: maximum_float(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-NEXT: .reg .b32 %f<4>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
@@ -1355,7 +1355,7 @@
;
; CHECK-SM80-NOF16-LABEL: maximum_float(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
@@ -1372,7 +1372,7 @@
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
@@ -1393,7 +1393,7 @@
;
; CHECK-F16-LABEL: maximum_float_ftz(
; CHECK-F16: {
-; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-NEXT: .reg .b32 %f<4>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
@@ -1404,7 +1404,7 @@
;
; CHECK-SM80-NOF16-LABEL: maximum_float_ftz(
; CHECK-SM80-NOF16: {
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
@@ -1421,7 +1421,7 @@
; CHECK: {
; CHECK-NEXT: .reg .pred %p<5>;
; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-NEXT: .reg .f64 %fd<8>;
+; CHECK-NEXT: .reg .b64 %fd<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [maximum_double_param_0];
@@ -1449,7 +1449,7 @@
; CHECK-NOF16-NEXT: .reg .pred %p<11>;
; CHECK-NOF16-NEXT: .reg .b16 %rs<15>;
; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_v2half_param_0];
@@ -1502,7 +1502,7 @@
; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;
; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<15>;
; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
-; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %f<7>;
; CHECK-SM80-NOF16-EMPTY:
; CHECK-SM80-NOF16-NEXT: // %bb.0:
; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_v2half_param_0];
@@ -1547,7 +1547,7 @@
define float @fma_float(float %a, float %b, float %c) {
; CHECK-LABEL: fma_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fma_float_param_0];
@@ -1563,7 +1563,7 @@
define float @fma_float_ftz(float %a, float %b, float %c) #1 {
; CHECK-LABEL: fma_float_ftz(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [fma_float_ftz_param_0];
@@ -1579,7 +1579,7 @@
define double @fma_double(double %a, double %b, double %c) {
; CHECK-LABEL: fma_double(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [fma_double_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/misched_func_call.ll b/llvm/test/CodeGen/NVPTX/misched_func_call.ll
index e0d0197..fb4c653 100644
--- a/llvm/test/CodeGen/NVPTX/misched_func_call.ll
+++ b/llvm/test/CodeGen/NVPTX/misched_func_call.ll
@@ -9,7 +9,7 @@
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<11>;
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: .reg .f64 %fd<6>;
+; CHECK-NEXT: .reg .b64 %fd<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %bb
; CHECK-NEXT: ld.param.u32 %r4, [wombat_param_2];
diff --git a/llvm/test/CodeGen/NVPTX/param-add.ll b/llvm/test/CodeGen/NVPTX/param-add.ll
index afabc11..c8daf3b 100644
--- a/llvm/test/CodeGen/NVPTX/param-add.ll
+++ b/llvm/test/CodeGen/NVPTX/param-add.ll
@@ -15,7 +15,7 @@
; CHECK-LABEL: test(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u8 %r1, [test_param_0+1];
diff --git a/llvm/test/CodeGen/NVPTX/rcp-opt.ll b/llvm/test/CodeGen/NVPTX/rcp-opt.ll
index 31fd8eb..0b020b7 100644
--- a/llvm/test/CodeGen/NVPTX/rcp-opt.ll
+++ b/llvm/test/CodeGen/NVPTX/rcp-opt.ll
@@ -9,7 +9,7 @@
define double @test1(double %in) {
; CHECK-LABEL: test1(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [test1_param_0];
@@ -27,7 +27,7 @@
define double @test2(double %in) {
; CHECK-LABEL: test2(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [test2_param_0];
@@ -44,7 +44,7 @@
define double @test3(double %in) {
; CHECK-LABEL: test3(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-NEXT: .reg .b64 %fd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [test3_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
index 2a12e9b..020a61a 100644
--- a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
@@ -115,7 +115,7 @@
define float @reduce_fadd_float(<8 x float> %in) {
; CHECK-LABEL: reduce_fadd_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<17>;
+; CHECK-NEXT: .reg .b32 %f<17>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fadd_float_param_0+16];
@@ -137,7 +137,7 @@
define float @reduce_fadd_float_reassoc(<8 x float> %in) {
; CHECK-LABEL: reduce_fadd_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<17>;
+; CHECK-NEXT: .reg .b32 %f<17>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fadd_float_reassoc_param_0+16];
@@ -159,7 +159,7 @@
define float @reduce_fadd_float_reassoc_nonpow2(<7 x float> %in) {
; CHECK-LABEL: reduce_fadd_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<15>;
+; CHECK-NEXT: .reg .b32 %f<15>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fadd_float_reassoc_nonpow2_param_0+24];
@@ -274,7 +274,7 @@
define float @reduce_fmul_float(<8 x float> %in) {
; CHECK-LABEL: reduce_fmul_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmul_float_param_0+16];
@@ -295,7 +295,7 @@
define float @reduce_fmul_float_reassoc(<8 x float> %in) {
; CHECK-LABEL: reduce_fmul_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmul_float_reassoc_param_0+16];
@@ -316,7 +316,7 @@
define float @reduce_fmul_float_reassoc_nonpow2(<7 x float> %in) {
; CHECK-LABEL: reduce_fmul_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<14>;
+; CHECK-NEXT: .reg .b32 %f<14>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fmul_float_reassoc_nonpow2_param_0+24];
@@ -404,7 +404,7 @@
;
; CHECK-LABEL: reduce_fmax_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmax_float_param_0+16];
@@ -426,7 +426,7 @@
;
; CHECK-LABEL: reduce_fmax_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmax_float_reassoc_param_0+16];
@@ -448,7 +448,7 @@
;
; CHECK-LABEL: reduce_fmax_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<14>;
+; CHECK-NEXT: .reg .b32 %f<14>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fmax_float_reassoc_nonpow2_param_0+24];
@@ -536,7 +536,7 @@
;
; CHECK-LABEL: reduce_fmin_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmin_float_param_0+16];
@@ -558,7 +558,7 @@
;
; CHECK-LABEL: reduce_fmin_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmin_float_reassoc_param_0+16];
@@ -580,7 +580,7 @@
;
; CHECK-LABEL: reduce_fmin_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<14>;
+; CHECK-NEXT: .reg .b32 %f<14>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fmin_float_reassoc_nonpow2_param_0+24];
@@ -668,7 +668,7 @@
;
; CHECK-LABEL: reduce_fmaximum_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmaximum_float_param_0+16];
@@ -690,7 +690,7 @@
;
; CHECK-LABEL: reduce_fmaximum_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fmaximum_float_reassoc_param_0+16];
@@ -712,7 +712,7 @@
;
; CHECK-LABEL: reduce_fmaximum_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<14>;
+; CHECK-NEXT: .reg .b32 %f<14>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fmaximum_float_reassoc_nonpow2_param_0+24];
@@ -800,7 +800,7 @@
;
; CHECK-LABEL: reduce_fminimum_float(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fminimum_float_param_0+16];
@@ -822,7 +822,7 @@
;
; CHECK-LABEL: reduce_fminimum_float_reassoc(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<16>;
+; CHECK-NEXT: .reg .b32 %f<16>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.f32 {%f5, %f6, %f7, %f8}, [reduce_fminimum_float_reassoc_param_0+16];
@@ -844,7 +844,7 @@
;
; CHECK-LABEL: reduce_fminimum_float_reassoc_nonpow2(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<14>;
+; CHECK-NEXT: .reg .b32 %f<14>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f7, [reduce_fminimum_float_reassoc_nonpow2_param_0+24];
diff --git a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
index af113e7..ed78529 100644
--- a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
+++ b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
@@ -7,7 +7,7 @@
; CHECK-LABEL: redux_sync_fmin(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_param_0];
@@ -24,7 +24,7 @@
; CHECK-LABEL: redux_sync_fmin_abs(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_abs_param_0];
@@ -41,7 +41,7 @@
; CHECK-LABEL: redux_sync_fmin_NaN(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_NaN_param_0];
@@ -58,7 +58,7 @@
; CHECK-LABEL: redux_sync_fmin_abs_NaN(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_abs_NaN_param_0];
@@ -75,7 +75,7 @@
; CHECK-LABEL: redux_sync_fmax(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_param_0];
@@ -92,7 +92,7 @@
; CHECK-LABEL: redux_sync_fmax_abs(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_abs_param_0];
@@ -109,7 +109,7 @@
; CHECK-LABEL: redux_sync_fmax_NaN(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_NaN_param_0];
@@ -126,7 +126,7 @@
; CHECK-LABEL: redux_sync_fmax_abs_NaN(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_abs_NaN_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/reg-types.ll b/llvm/test/CodeGen/NVPTX/reg-types.ll
index cf2433a..7b4ebca 100644
--- a/llvm/test/CodeGen/NVPTX/reg-types.ll
+++ b/llvm/test/CodeGen/NVPTX/reg-types.ll
@@ -25,9 +25,9 @@
%u64 = alloca i64, align 8
; CHECK-DAG: .reg .b64 %rd<
%f32 = alloca float, align 4
-; CHECK-DAG: .reg .f32 %f<
+; CHECK-DAG: .reg .b32 %f<
%f64 = alloca double, align 8
-; CHECK-DAG: .reg .f64 %fd<
+; CHECK-DAG: .reg .b64 %fd<
; Verify that we use correct register types.
store i8 1, ptr %s8, align 1
diff --git a/llvm/test/CodeGen/NVPTX/st-param-imm.ll b/llvm/test/CodeGen/NVPTX/st-param-imm.ll
index ab14476..e8ad689 100644
--- a/llvm/test/CodeGen/NVPTX/st-param-imm.ll
+++ b/llvm/test/CodeGen/NVPTX/st-param-imm.ll
@@ -403,7 +403,7 @@
define void @st_param_v2_f32_ir(float %val) {
; CHECK-LABEL: st_param_v2_f32_ir(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v2_f32_ir_param_0];
@@ -425,7 +425,7 @@
define void @st_param_v2_f32_ri(float %val) {
; CHECK-LABEL: st_param_v2_f32_ri(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v2_f32_ri_param_0];
@@ -467,7 +467,7 @@
define void @st_param_v2_f64_ir(double %val) {
; CHECK-LABEL: st_param_v2_f64_ir(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [st_param_v2_f64_ir_param_0];
@@ -489,7 +489,7 @@
define void @st_param_v2_f64_ri(double %val) {
; CHECK-LABEL: st_param_v2_f64_ri(
; CHECK: {
-; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-NEXT: .reg .b64 %fd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f64 %fd1, [st_param_v2_f64_ri_param_0];
@@ -1648,7 +1648,7 @@
define void @st_param_v4_f32_irrr(float %b, float %c, float %d) {
; CHECK-LABEL: st_param_v4_f32_irrr(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_irrr_param_0];
@@ -1674,7 +1674,7 @@
define void @st_param_v4_f32_rirr(float %a, float %c, float %d) {
; CHECK-LABEL: st_param_v4_f32_rirr(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_rirr_param_0];
@@ -1700,7 +1700,7 @@
define void @st_param_v4_f32_rrir(float %a, float %b, float %d) {
; CHECK-LABEL: st_param_v4_f32_rrir(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_rrir_param_0];
@@ -1726,7 +1726,7 @@
define void @st_param_v4_f32_rrri(float %a, float %b, float %c) {
; CHECK-LABEL: st_param_v4_f32_rrri(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-NEXT: .reg .b32 %f<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_rrri_param_0];
@@ -1752,7 +1752,7 @@
define void @st_param_v4_f32_iirr(float %c, float %d) {
; CHECK-LABEL: st_param_v4_f32_iirr(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_iirr_param_0];
@@ -1777,7 +1777,7 @@
define void @st_param_v4_f32_irir(float %b, float %d) {
; CHECK-LABEL: st_param_v4_f32_irir(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_irir_param_0];
@@ -1802,7 +1802,7 @@
define void @st_param_v4_f32_irri(float %b, float %c) {
; CHECK-LABEL: st_param_v4_f32_irri(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_irri_param_0];
@@ -1827,7 +1827,7 @@
define void @st_param_v4_f32_riir(float %a, float %d) {
; CHECK-LABEL: st_param_v4_f32_riir(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_riir_param_0];
@@ -1852,7 +1852,7 @@
define void @st_param_v4_f32_riri(float %a, float %c) {
; CHECK-LABEL: st_param_v4_f32_riri(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_riri_param_0];
@@ -1877,7 +1877,7 @@
define void @st_param_v4_f32_rrii(float %a, float %b) {
; CHECK-LABEL: st_param_v4_f32_rrii(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-NEXT: .reg .b32 %f<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_rrii_param_0];
@@ -1902,7 +1902,7 @@
define void @st_param_v4_f32_iiir(float %d) {
; CHECK-LABEL: st_param_v4_f32_iiir(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_iiir_param_0];
@@ -1926,7 +1926,7 @@
define void @st_param_v4_f32_iiri(float %c) {
; CHECK-LABEL: st_param_v4_f32_iiri(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_iiri_param_0];
@@ -1950,7 +1950,7 @@
define void @st_param_v4_f32_irii(float %b) {
; CHECK-LABEL: st_param_v4_f32_irii(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_irii_param_0];
@@ -1974,7 +1974,7 @@
define void @st_param_v4_f32_riii(float %a) {
; CHECK-LABEL: st_param_v4_f32_riii(
; CHECK: {
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.f32 %f1, [st_param_v4_f32_riii_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll b/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll
index 7a7904a..3afff32 100644
--- a/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll
+++ b/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll
@@ -14,7 +14,7 @@
; CHECK-LABEL: foo(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -38,7 +38,7 @@
; CHECK-LABEL: bar(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b32 %f<2>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
diff --git a/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll b/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll
index 61837bd..4e4e3f3 100644
--- a/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll
+++ b/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll
@@ -14,7 +14,7 @@
; CHECK-LABEL: foo(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -38,7 +38,7 @@
; CHECK-LABEL: bar(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-NEXT: .reg .b32 %f<5>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
@@ -61,7 +61,7 @@
; CHECK-LABEL: baz(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .f32 %f<8>;
+; CHECK-NEXT: .reg .b32 %f<8>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index 35db489..9da3614 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -13,7 +13,7 @@
; CHECK-PTX: {
; CHECK-PTX-NEXT: .reg .b32 %r<11>;
; CHECK-PTX-NEXT: .reg .b64 %rd<11>;
-; CHECK-PTX-NEXT: .reg .f64 %fd<7>;
+; CHECK-PTX-NEXT: .reg .b64 %fd<7>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: ld.param.u32 %r1, [variadics1_param_0];
diff --git a/llvm/test/DebugInfo/NVPTX/debug-info.ll b/llvm/test/DebugInfo/NVPTX/debug-info.ll
index fa2925a..1fc945b 100644
--- a/llvm/test/DebugInfo/NVPTX/debug-info.ll
+++ b/llvm/test/DebugInfo/NVPTX/debug-info.ll
@@ -20,7 +20,7 @@
; CHECK: )
; CHECK: {
; CHECK-DAG: .reg .pred %p<2>;
-; CHECK-DAG: .reg .f32 %f<5>;
+; CHECK-DAG: .reg .b32 %f<5>;
; CHECK-DAG: .reg .b32 %r<6>;
; CHECK-DAG: .reg .b64 %rd<8>;
; CHECK: .loc [[DEBUG_INFO_CU:[0-9]+]] 5 0
