diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp index 142388893082a..6b9797c3e6aae 100644 --- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp @@ -26,9 +26,9 @@ using namespace llvm; 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 e46657e4a582f..8f0964c2d5eba 100644 --- a/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll +++ b/llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll @@ -45,7 +45,7 @@ define half @fh(ptr %p) { ; 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 @@ define half @fh(ptr %p) { ; 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 half @fh(ptr %p) { 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 @@ define float @ff(ptr %p) { ; ; 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 6c3514c1ad946..5949de335b8cf 100644 --- a/llvm/test/CodeGen/NVPTX/and-or-setcc.ll +++ b/llvm/test/CodeGen/NVPTX/and-or-setcc.ll @@ -9,7 +9,7 @@ define i1 @and_ord(float %a, float %b) { ; 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 @@ define i1 @or_uno(float %a, float %b) { ; 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 bb04aa856d656..16de80d55a054 100644 --- a/llvm/test/CodeGen/NVPTX/atomics.ll +++ b/llvm/test/CodeGen/NVPTX/atomics.ll @@ -351,7 +351,7 @@ declare float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val) 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 @@ declare float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %v 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 @@ declare float @llvm.nvvm.atomic.load.add.f32.p3(ptr addrspace(3) %addr, float %v 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 @atomic_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) { 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 @@ define half @atomicrmw_add_f16_generic(ptr %addr, half %val) { ; 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 half @atomicrmw_add_f16_generic(ptr %addr, half %val) { 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_addrspace1(ptr addrspace(1) %addr, float %val) { 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 b97cb6fa3cbe4..6be13c3a6fdec 100644 --- a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll @@ -19,7 +19,7 @@ define bfloat @test_fadd(bfloat %0, bfloat %1) { ; 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 @@ define bfloat @test_fadd(bfloat %0, bfloat %1) { ; 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 @@ define bfloat @test_fsub(bfloat %0, bfloat %1) { ; 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 @@ define bfloat @test_fsub(bfloat %0, bfloat %1) { ; 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 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define float @test_fpext_float(bfloat %a) #0 { ; 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 @@ define float @test_fpext_float(bfloat %a) #0 { ; 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 @@ define float @test_fpext_float(bfloat %a) #0 { ; 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 @@ define float @test_fpext_float(bfloat %a) #0 { ; 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 @@ define bfloat @test_fptrunc_float(float %a) #0 { ; 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 @@ define bfloat @test_fptrunc_float(float %a) #0 { ; 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 @@ define bfloat @test_fptrunc_float(float %a) #0 { ; 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 @@ define bfloat @test_fptrunc_float(float %a) #0 { ; 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 @@ define bfloat @test_fadd_imm_1(bfloat %a) #0 { ; 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 @@ define bfloat @test_fadd_imm_1(bfloat %a) #0 { ; 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 @@ define bfloat @test_select_cc_bf16_f64(double %a, double %b, bfloat %c, bfloat % ; 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 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 { ; 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 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 { ; 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 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 { ; 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 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 { ; 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 @@ define i16 @test_fptosi_i16(bfloat %a) { ; 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 @@ define i16 @test_fptosi_i16(bfloat %a) { ; 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 @@ define i16 @test_fptosi_i16(bfloat %a) { ; 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 @@ define i16 @test_fptoui_i16(bfloat %a) { ; 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 @@ define i16 @test_fptoui_i16(bfloat %a) { ; 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 @@ define i16 @test_fptoui_i16(bfloat %a) { ; 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 @@ define bfloat @test_sitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_sitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_sitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_uitofp_i8(i8 %a) { ; 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 @@ define bfloat @test_uitofp_i8(i8 %a) { ; 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 @@ define bfloat @test_uitofp_i8(i8 %a) { ; 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 @@ define bfloat @test_uitofp_i1(i1 %a) { ; 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 @@ define bfloat @test_uitofp_i1(i1 %a) { ; 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 @@ define bfloat @test_uitofp_i1(i1 %a) { ; 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 @@ define bfloat @test_uitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_uitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_uitofp_i16(i16 %a) { ; 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 @@ define bfloat @test_uitofp_i32(i32 %a) { ; 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 @@ define bfloat @test_uitofp_i32(i32 %a) { ; 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 @@ define bfloat @test_uitofp_i32(i32 %a) { ; 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 @@ define bfloat @test_uitofp_i64(i64 %a) { ; 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 @@ define bfloat @test_uitofp_i64(i64 %a) { ; 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 @@ define bfloat @test_uitofp_i64(i64 %a) { ; 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 @@ define bfloat @test_roundeven(bfloat %a) { ; 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 @@ define bfloat @test_roundeven(bfloat %a) { ; 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 @@ define bfloat @test_roundeven(bfloat %a) { ; 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 @@ define bfloat @test_maximum(bfloat %a, bfloat %b) { ; 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 @@ define bfloat @test_maxnum(bfloat %a, bfloat %b) { ; 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 @@ define <2 x bfloat> @test_maximum_v2(<2 x bfloat> %a, <2 x bfloat> %b) { ; 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 @@ define <2 x bfloat> @test_maxnum_v2(<2 x bfloat> %a, <2 x bfloat> %b) { ; 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 @@ define <2 x bfloat> @test_maxnum_v2(<2 x bfloat> %a, <2 x bfloat> %b) { } 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 fdf481e1767a9..5ab684adac58e 100644 --- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll +++ b/llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll @@ -12,7 +12,7 @@ define <2 x bfloat> @test_sin(<2 x bfloat> %a) #0 #1 { ; 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 @@ define <2 x bfloat> @test_cos(<2 x bfloat> %a) #0 #1 { ; 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 706a88f2cb901..677f0d795dde8 100644 --- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll @@ -131,7 +131,7 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 { ; 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 @@ define <2 x bfloat> @test_select_cc(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloa ; 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 @@ define <2 x float> @test_select_cc_f32_bf16(<2 x float> %a, <2 x float> %b, ; 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 @@ define <2 x float> @test_select_cc_f32_bf16(<2 x float> %a, <2 x float> %b, ; 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 @@ define <2 x bfloat> @test_select_cc_bf16_f32(<2 x bfloat> %a, <2 x bfloat> %b, ; 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 @@ define <2 x bfloat> @test_fptrunc_2xfloat(<2 x float> %a) #0 { ; 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 @@ define <2 x float> @test_fpext_2xfloat(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_sqrt(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_floor(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_ceil(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_trunc(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_rint(<2 x bfloat> %a) #0 { ; 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 @@ define <2 x bfloat> @test_round(<2 x bfloat> %a) #0 { ; 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 93da39137afd8..670e112c26c76 100644 --- a/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll +++ b/llvm/test/CodeGen/NVPTX/convert-fp-i8.ll @@ -8,7 +8,7 @@ define i8 @cvt_u8_f32(float %x) { ; 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 @@ define i8 @cvt_u8_f64(double %x) { ; 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 @@ define float @cvt_f32_i8(i8 %x) { ; 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 @@ define double @cvt_f64_i8(i8 %x) { ; 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 @@ define float @cvt_f32_s8(i8 %x) { ; 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 @@ define double @cvt_f64_s8(i8 %x) { ; 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 @@ define i8 @cvt_s8_f32(float %x) { ; 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 @@ define i8 @cvt_s8_f64(double %x) { ; 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 f92822f7e0c16..7230872b3427c 100644 --- a/llvm/test/CodeGen/NVPTX/convert-sm100.ll +++ b/llvm/test/CodeGen/NVPTX/convert-sm100.ll @@ -11,7 +11,7 @@ define i32 @cvt_rn_satf_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rn_relu_satf_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rz_satf_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rz_relu_satf_tf32_f32(float %f1) { ; 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 f0dd5f084026b..04d7a65f9e40e 100644 --- a/llvm/test/CodeGen/NVPTX/convert-sm100a.ll +++ b/llvm/test/CodeGen/NVPTX/convert-sm100a.ll @@ -11,7 +11,7 @@ define i16 @cvt_rn_sf_e2m3x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rn_relu_sf_e2m3x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rn_sf_e3m2x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rn_relu_sf_e3m2x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rz_ue8m0x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rz_sf_ue8m0x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rp_ue8m0x2_f32(float %f1, float %f2) { ; 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 @@ define i16 @cvt_rp_sf_ue8m0x2_f32(float %f1, float %f2) { ; 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 aebc28b1cfea3..eb7a6bdd222bb 100644 --- a/llvm/test/CodeGen/NVPTX/convert-sm80.ll +++ b/llvm/test/CodeGen/NVPTX/convert-sm80.ll @@ -7,7 +7,7 @@ define <2 x bfloat> @cvt_rn_bf16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x bfloat> @cvt_rn_relu_bf16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x bfloat> @cvt_rz_bf16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x bfloat> @cvt_rz_relu_bf16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x half> @cvt_rn_f16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x half> @cvt_rn_relu_f16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x half> @cvt_rz_f16x2_f32(float %f1, float %f2) { ; 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 @@ define <2 x half> @cvt_rz_relu_f16x2_f32(float %f1, float %f2) { ; 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 @@ define bfloat @cvt_rn_bf16_f32(float %f1) { ; 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 @@ define bfloat @cvt_rn_relu_bf16_f32(float %f1) { ; 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 @@ define bfloat @cvt_rz_bf16_f32(float %f1) { ; 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 @@ define bfloat @cvt_rz_relu_bf16_f32(float %f1) { ; 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 @@ define i32 @cvt_rna_tf32_f32(float %f1) { ; 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 @@ define <2 x bfloat> @fold_ff2bf16x2(float %lo, float %hi) { ; 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 @@ define <2 x half> @fold_ff2f16x2(float %lo, float %hi) { ; 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 5f610e0e91f88..340117f98cd94 100644 --- a/llvm/test/CodeGen/NVPTX/convert-sm90.ll +++ b/llvm/test/CodeGen/NVPTX/convert-sm90.ll @@ -11,7 +11,7 @@ define i32 @cvt_rn_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rn_relu_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rz_tf32_f32(float %f1) { ; 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 @@ define i32 @cvt_rz_relu_tf32_f32(float %f1) { ; 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 843ef4dbde367..2e305e683d777 100644 --- a/llvm/test/CodeGen/NVPTX/copysign.ll +++ b/llvm/test/CodeGen/NVPTX/copysign.ll @@ -8,7 +8,7 @@ target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f3 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 float @fcopysign_f_f(float %a, float %b) { 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 @@ define float @fcopysign_f_d(float %a, double %b) { ; 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 @@ define float @fcopysign_f_h(float %a, half %b) { ; 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 @@ define double @fcopysign_d_f(double %a, float %b) { ; 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 @@ define double @fcopysign_d_h(double %a, half %b) { ; 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 92b73799d6f1d..a233616563085 100644 --- a/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll +++ b/llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll @@ -59,9 +59,9 @@ define void @test_distributed_shared_cluster_float_atomic(ptr addrspace(7) %dsme ; 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 4f9d58758ca9e..f8711e3a83591 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 d12653e813bd1..d3aaedf84bce9 100644 --- a/llvm/test/CodeGen/NVPTX/f16-abs.ll +++ b/llvm/test/CodeGen/NVPTX/f16-abs.ll @@ -49,7 +49,7 @@ define half @test_fabs(half %a) { ; 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 e9edabd1ee8af..e854e5a6e5aaa 100644 --- a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll @@ -109,7 +109,7 @@ define <2 x half> @test_fadd(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_fadd_imm_0(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_fadd_imm_1(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_fsub(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_fneg(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_fmul(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_fdiv(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_frem(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, < ; 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 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b, ; 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 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b, ; 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 @@ define <2 x half> @test_select_cc_f16_f32(<2 x half> %a, <2 x half> %b, ; 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 @@ define <2 x i1> @test_fcmp_une(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ueq(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ugt(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_uge(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ult(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ule(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_uno(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_one(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_oeq(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ogt(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_oge(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_olt(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ole(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x i1> @test_fcmp_ord(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_uitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_sitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_fptrunc_2xfloat(<2 x float> %a) #0 { ; 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 @@ define <2 x half> @test_fptrunc_2xdouble(<2 x double> %a) #0 { ; 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 @@ define <2 x float> @test_fpext_2xfloat(<2 x half> %a) #0 { ; 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 @@ define <2 x double> @test_fpext_2xdouble(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_bitcast_float_to_2xhalf(float %a) #0 { ; 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 @@ define float @test_bitcast_2xhalf_to_float(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_sqrt(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_sin(<2 x half> %a) #0 #1 { ; 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 @@ define <2 x half> @test_cos(<2 x half> %a) #0 #1 { ; 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 @@ define <2 x half> @test_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 { ; 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 @@ define <2 x half> @test_fabs(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_minnum(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_maxnum(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 { ; 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 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 { ; 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 @@ define <2 x half> @test_copysign_f64(<2 x half> %a, <2 x double> %b) #0 { ; 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 @@ define <2 x half> @test_copysign_f64(<2 x half> %a, <2 x double> %b) #0 { ; 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 @@ define <2 x float> @test_copysign_extended(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x float> @test_copysign_extended(<2 x half> %a, <2 x half> %b) #0 { ; 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 @@ define <2 x half> @test_round(<2 x half> %a) #0 { ; 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 @@ define <2 x half> @test_fmuladd(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 ; 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 c9eff2a8ff17d..2c5c8146fbf61 100644 --- a/llvm/test/CodeGen/NVPTX/f32-ex2.ll +++ b/llvm/test/CodeGen/NVPTX/f32-ex2.ll @@ -9,7 +9,7 @@ declare float @llvm.nvvm.ex2.approx.f(float) 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(float %0) { 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 43c521978fed8..9dac3083d6cb8 100644 --- a/llvm/test/CodeGen/NVPTX/f32-lg2.ll +++ b/llvm/test/CodeGen/NVPTX/f32-lg2.ll @@ -10,7 +10,7 @@ declare float @llvm.nvvm.lg2.approx.ftz.f(float) 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(float %0) { 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 dd9ef220a9b47..d9c5a527b901c 100644 --- a/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll +++ b/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll @@ -18,7 +18,7 @@ declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat>) 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(float %a) { 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 float @fabs_float_ftz(float %a) { 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 7e485dca65764..4664d700209fa 100644 --- a/llvm/test/CodeGen/NVPTX/fexp2.ll +++ b/llvm/test/CodeGen/NVPTX/fexp2.ll @@ -13,7 +13,7 @@ target triple = "nvptx64-nvidia-cuda" 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 @@ define float @exp2_test(float %in) { ; ; 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 @@ define float @exp2_test(float %in) { ; ; 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 @@ entry: 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 @@ define float @exp2_ftz_test(float %in) #0 { ; ; 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 @@ define float @exp2_ftz_test(float %in) #0 { ; ; 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 @@ entry: 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 @@ define <2 x float> @exp2_test_v(<2 x float> %in) { ; ; 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 @@ define <2 x float> @exp2_test_v(<2 x float> %in) { ; ; 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 @@ define half @exp2_f16_test(half %in) { ; 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 @@ define half @exp2_f16_ftz_test(half %in) #0 { ; 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 @@ define <2 x half> @exp2_f16_test_v(<2 x half> %in) { ; 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 @@ define bfloat @exp2_bf16_test(bfloat %in) { ; 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 @@ define bfloat @exp2_bf16_test(bfloat %in) { ; 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 @@ define <2 x bfloat> @exp2_bf16_test_v(<2 x bfloat> %in) { ; 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 @@ define <2 x bfloat> @exp2_bf16_test_v(<2 x bfloat> %in) { ; 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 ff762dcf74b2f..4dfed3dd944ae 100644 --- a/llvm/test/CodeGen/NVPTX/flog2.ll +++ b/llvm/test/CodeGen/NVPTX/flog2.ll @@ -7,7 +7,7 @@ target triple = "nvptx64-nvidia-cuda" 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 @@ entry: 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 @@ entry: 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 @@ define half @log2_f16_test(half %in) { ; 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 @@ define half @log2_f16_ftz_test(half %in) #0 { ; 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 @@ define <2 x half> @log2_f16_test_v(<2 x half> %in) { ; 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 @@ define bfloat @log2_bf16_test(bfloat %in) { ; 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 @@ define bfloat @log2_bf16_ftz_test(bfloat %in) #0 { ; 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 @@ define <2 x bfloat> @log2_bf16_test_v(<2 x bfloat> %in) { ; 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 7dce894620e6b..9051a0bce14cd 100644 --- a/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll +++ b/llvm/test/CodeGen/NVPTX/fma-relu-contract.ll @@ -198,7 +198,7 @@ define half @fma_f16_expanded_maxnum_no_nans(half %a, half %b, half %c) #0 { ; 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 @@ define bfloat @fma_bf16_expanded_unsafe_with_nans(bfloat %a, bfloat %b, bfloat % ; 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 @@ define bfloat @fma_bf16_expanded_no_nans(bfloat %a, bfloat %b, bfloat %c) #0 { ; 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 @@ define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat ; 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 @@ define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat ; 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 @@ define bfloat @fma_bf16_expanded_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) ; 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 @@ define <2 x half> @fma_f16x2_expanded_maxnum_no_nans(<2 x half> %a, <2 x half> % ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_unsafe_with_nans(<2 x bfloat> %a, <2 x ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans(<2 x bfloat> %a, <2 x bfloat> % ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloa ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloa ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_maxnum_no_nans(<2 x bfloat> %a, <2 x bf ; 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 eb51d7db81372..73f808f1e06ee 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 @@ define half @fma_f16_maxnum_no_nans(half %a, half %b, half %c) #0 { ; 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 @@ define bfloat @fma_bf16_no_nans(bfloat %a, bfloat %b, bfloat %c) #0 { ; 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 @@ define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloa ; 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 @@ define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloa ; 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 @@ define bfloat @fma_bf16_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) #0 { ; 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 @@ define <2 x half> @fma_f16x2_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x h ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x b ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 ; 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 @@ define <2 x bfloat> @fma_bf16x2_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, ; 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 a3545f5171425..b94fa5a24b502 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 @@ define half @fma_f16_expanded_maxnum_no_nans(half %a, half %b, half %c) { ; 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 @@ define bfloat @fma_bf16_expanded_no_nans(bfloat %a, bfloat %b, bfloat %c) { ; 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 @@ define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat ; 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 @@ define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat ; 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 @@ define bfloat @fma_bf16_expanded_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) ; 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 @@ define <2 x half> @fma_f16x2_expanded_maxnum_no_nans(<2 x half> %a, <2 x half> % ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans(<2 x bfloat> %a, <2 x bfloat> % ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloa ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloa ; 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 @@ define <2 x bfloat> @fma_bf16x2_expanded_maxnum_no_nans(<2 x bfloat> %a, <2 x bf ; 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 @@ define half @fma_f16_maxnum_no_nans(half %a, half %b, half %c) { ; 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 @@ define bfloat @fma_bf16_no_nans(bfloat %a, bfloat %b, bfloat %c) { ; 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 @@ define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloa ; 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 @@ define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloa ; 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 @@ define bfloat @fma_bf16_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) { ; 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 @@ define <2 x half> @fma_f16x2_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x h ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x b ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 ; 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 @@ define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 ; 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 @@ define <2 x bfloat> @fma_bf16x2_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, ; 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 ea5da6ee57f65..bd559ea157feb 100644 --- a/llvm/test/CodeGen/NVPTX/fp-contract.ll +++ b/llvm/test/CodeGen/NVPTX/fp-contract.ll @@ -15,7 +15,7 @@ target triple = "nvptx64-unknown-cuda" 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 @@ define float @t0(float %a, float %b, float %c) { ; ; 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 @t0(float %a, float %b, float %c) { 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 @@ define float @t1(float %a, float %b) { ; ; 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 @t1(float %a, float %b) { 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 @t2(float %a, float %b) { 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 73debfbfcdf49..4077f6d1eb21b 100644 --- a/llvm/test/CodeGen/NVPTX/frem.ll +++ b/llvm/test/CodeGen/NVPTX/frem.ll @@ -9,7 +9,7 @@ define half @frem_f16(half %a, half %b) { ; 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 @@ define half @frem_f16(half %a, half %b) { ; 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 half @frem_f16(half %a, half %b) { 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 @@ define float @frem_f32(float %a, float %b) { ; 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 float @frem_f32(float %a, float %b) { 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 @@ define double @frem_f64(double %a, double %b) { ; 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 @@ define half @frem_f16_ninf(half %a, half %b) { ; 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 @@ define half @frem_f16_ninf(half %a, half %b) { ; 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 half @frem_f16_ninf(half %a, half %b) { 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 @@ define float @frem_f32_ninf(float %a, float %b) { ; ; 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 float @frem_f32_ninf(float %a, float %b) { 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 @@ define double @frem_f64_ninf(double %a, double %b) { ; ; 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 double @frem_f64_ninf(double %a, double %b) { 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 @@ define float @frem_f32_imm1(float %a) { ; ; 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_imm1(float %a) { 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 @@ define float @frem_f32_imm2(float %a) { ; 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 1b779dbe5abeb..65edcf2e07159 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -1141,7 +1141,7 @@ define <4 x i8> @test_bitcast_float_to_4xi8(float %a) #0 { ; 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 @@ define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 { ; 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 cc6af060d6c0a..01c51bb72d055 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 float @test_fabsf(float %f) { 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 double @test_fabs(double %d) { 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_nvvm_sqrt(float %a) { 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 16a0189e784bd..2fe2d28320f06 100644 --- a/llvm/test/CodeGen/NVPTX/ldg-invariant.ll +++ b/llvm/test/CodeGen/NVPTX/ldg-invariant.ll @@ -27,7 +27,7 @@ define half @ld_global_v2f16(ptr addrspace(1) %ptr) { ; 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 @@ define half @ld_global_v4f16(ptr addrspace(1) %ptr) { ; 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 @@ define half @ld_global_v8f16(ptr addrspace(1) %ptr) { ; 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 4c5c44a9bf44d..2c1550aa082f0 100644 --- a/llvm/test/CodeGen/NVPTX/ldu-ldg.ll +++ b/llvm/test/CodeGen/NVPTX/ldu-ldg.ll @@ -104,7 +104,7 @@ define ptr @test_ldu_p(ptr addrspace(1) %ptr) { 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 @@ define double @test_ldu_f64(ptr addrspace(1) %ptr) { ; 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 ptr @test_ldg_p(ptr addrspace(1) %ptr) { 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 @@ define double @test_ldg_f64(ptr addrspace(1) %ptr) { ; 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 ed94cb416f472..cb2e247bd78c1 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_i64(ptr %a) { 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 @@ define void @generic_double(ptr %a) { ; 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_i64(ptr %a) { 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 @@ define void @generic_volatile_double(ptr %a) { ; 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_i64(ptr %a) { 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 @@ define void @generic_unordered_sys_float(ptr %a) { ; ; 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 @@ define void @generic_unordered_sys_double(ptr %a) { ; 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 @@ define void @generic_unordered_sys_double(ptr %a) { ; 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_i64(ptr %a) { 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 @@ define void @generic_unordered_volatile_sys_double(ptr %a) { ; 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_i64(ptr %a) { 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 @@ define void @generic_monotonic_sys_float(ptr %a) { ; ; 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 @@ define void @generic_monotonic_sys_double(ptr %a) { ; 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 @@ define void @generic_monotonic_sys_double(ptr %a) { ; 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_i64(ptr %a) { 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 @@ define void @generic_monotonic_volatile_sys_double(ptr %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_volatile_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_unordered_sys_float(ptr addrspace(1) %a) { ; ; 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 @@ define void @global_unordered_sys_double(ptr addrspace(1) %a) { ; 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 @@ define void @global_unordered_sys_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_unordered_volatile_sys_float(ptr addrspace(1) %a) { ; ; 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 @@ define void @global_unordered_volatile_sys_double(ptr addrspace(1) %a) { ; 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 @@ define void @global_unordered_volatile_sys_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_monotonic_sys_float(ptr addrspace(1) %a) { ; ; 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 @@ define void @global_monotonic_sys_double(ptr addrspace(1) %a) { ; 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 @@ define void @global_monotonic_sys_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(1) %a) { 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 @@ define void @global_monotonic_volatile_sys_float(ptr addrspace(1) %a) { ; ; 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 @@ define void @global_monotonic_volatile_sys_double(ptr addrspace(1) %a) { ; 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 @@ define void @global_monotonic_volatile_sys_double(ptr addrspace(1) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_volatile_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_unordered_sys_float(ptr addrspace(3) %a) { ; ; 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 @@ define void @shared_unordered_sys_double(ptr addrspace(3) %a) { ; 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 @@ define void @shared_unordered_sys_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_unordered_volatile_sys_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_monotonic_sys_float(ptr addrspace(3) %a) { ; ; 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 @@ define void @shared_monotonic_sys_double(ptr addrspace(3) %a) { ; 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 @@ define void @shared_monotonic_sys_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(3) %a) { 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 @@ define void @shared_monotonic_volatile_sys_double(ptr addrspace(3) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_double(ptr addrspace(5) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_volatile_double(ptr addrspace(5) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_unordered_sys_double(ptr addrspace(5) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_unordered_volatile_sys_double(ptr addrspace(5) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_monotonic_sys_double(ptr addrspace(5) %a) { ; 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_i64(ptr addrspace(5) %a) { 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 @@ define void @local_monotonic_volatile_sys_double(ptr addrspace(5) %a) { ; 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 ba397dca68f1b..3215fce964005 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_2xi64(ptr %a) { 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_2xfloat(ptr %a) { 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 @@ define void @generic_2xdouble(ptr %a) { ; 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_2xi64(ptr %a) { 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_2xfloat(ptr %a) { 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 @@ define void @generic_volatile_2xdouble(ptr %a) { ; 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_2xi64(ptr addrspace(1) %a) { 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_2xfloat(ptr addrspace(1) %a) { 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 @@ define void @global_2xdouble(ptr addrspace(1) %a) { ; 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_2xi64(ptr addrspace(1) %a) { 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_2xfloat(ptr addrspace(1) %a) { 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 @@ define void @global_volatile_2xdouble(ptr addrspace(1) %a) { ; 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_2xi64(ptr addrspace(3) %a) { 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_2xfloat(ptr addrspace(3) %a) { 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 @@ define void @shared_2xdouble(ptr addrspace(3) %a) { ; 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_2xi64(ptr addrspace(3) %a) { 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_2xfloat(ptr addrspace(3) %a) { 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 @@ define void @shared_volatile_2xdouble(ptr addrspace(3) %a) { ; 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_2xi64(ptr addrspace(5) %a) { 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_2xfloat(ptr addrspace(5) %a) { 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 @@ define void @local_2xdouble(ptr addrspace(5) %a) { ; 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_2xi64(ptr addrspace(5) %a) { 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_2xfloat(ptr addrspace(5) %a) { 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 @@ define void @local_volatile_2xdouble(ptr addrspace(5) %a) { ; 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 189f3421cd03a..a6d01c16c0ab8 100644 --- a/llvm/test/CodeGen/NVPTX/math-intrins.ll +++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll @@ -50,7 +50,7 @@ declare double @llvm.fma.f64(double, double, double) #0 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(float %a) { 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 float @ceil_float_ftz(float %a) #1 { 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 double @ceil_double(double %a) { 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(float %a) { 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 float @floor_float_ftz(float %a) #1 { 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 @@ define float @round_float(float %a) { ; 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 @@ define float @round_float_ftz(float %a) #1 { ; 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 @@ define double @round_double(double %a) { ; 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 double @round_double(double %a) { 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(float %a) { 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 float @nearbyint_float_ftz(float %a) #1 { 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 double @nearbyint_double(double %a) { 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(float %a) { 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 float @rint_float_ftz(float %a) #1 { 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 double @rint_double(double %a) { 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(float %a) { 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 float @roundeven_float_ftz(float %a) #1 { 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 double @roundeven_double(double %a) { 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(float %a) { 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 float @trunc_float_ftz(float %a) #1 { 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 double @trunc_double(double %a) { 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(float %a) { 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 float @abs_float_ftz(float %a) #1 { 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 @@ define half @minnum_half(half %a, half %b) { ; 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 @@ define half @minnum_half(half %a, half %b) { ; 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 half @minnum_half(half %a, half %b) { 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_float(float %a, float %b) { 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_imm1(float %a) { 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_imm2(float %a) { 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 float @minnum_float_ftz(float %a, float %b) #1 { 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 @@ define <2 x half> @minnum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define <2 x half> @minnum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define half @minimum_half(half %a, half %b) { ; 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 @@ define half @minimum_half(half %a, half %b) { ; 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 @@ define float @minimum_float(float %a, float %b) { ; 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 @@ define float @minimum_float(float %a, float %b) { ; ; 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 @@ define float @minimum_float(float %a, float %b) { ; ; 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 @@ define float @minimum_imm1(float %a) { ; 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 @@ define float @minimum_imm1(float %a) { ; ; 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 @@ define float @minimum_imm1(float %a) { ; ; 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 @@ define float @minimum_imm2(float %a) { ; 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 @@ define float @minimum_imm2(float %a) { ; ; 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 @@ define float @minimum_imm2(float %a) { ; ; 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 @@ define float @minimum_float_ftz(float %a, float %b) #1 { ; 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 @@ define float @minimum_float_ftz(float %a, float %b) #1 { ; ; 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 @@ define float @minimum_float_ftz(float %a, float %b) #1 { ; ; 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 @@ define double @minimum_double(double %a, double %b) { ; 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 @@ define <2 x half> @minimum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define <2 x half> @minimum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define half @maxnum_half(half %a, half %b) { ; 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 @@ define half @maxnum_half(half %a, half %b) { ; 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 half @maxnum_half(half %a, half %b) { 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_imm1(float %a) { 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_imm2(float %a) { 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(float %a, float %b) { 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 float @maxnum_float_ftz(float %a, float %b) #1 { 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 @@ define <2 x half> @maxnum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define <2 x half> @maxnum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define half @maximum_half(half %a, half %b) { ; 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 @@ define half @maximum_half(half %a, half %b) { ; 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 @@ define float @maximum_imm1(float %a) { ; 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 @@ define float @maximum_imm1(float %a) { ; ; 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 @@ define float @maximum_imm1(float %a) { ; ; 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 @@ define float @maximum_imm2(float %a) { ; 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 @@ define float @maximum_imm2(float %a) { ; ; 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 @@ define float @maximum_imm2(float %a) { ; ; 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 @@ define float @maximum_float(float %a, float %b) { ; 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 @@ define float @maximum_float(float %a, float %b) { ; ; 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 @@ define float @maximum_float(float %a, float %b) { ; ; 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 @@ define float @maximum_float_ftz(float %a, float %b) #1 { ; 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 @@ define float @maximum_float_ftz(float %a, float %b) #1 { ; ; 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 @@ define float @maximum_float_ftz(float %a, float %b) #1 { ; ; 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 @@ define double @maximum_double(double %a, double %b) { ; 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 @@ define <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 @@ define <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) { ; 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 <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) { 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(float %a, float %b, float %c) { 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 float @fma_float_ftz(float %a, float %b, float %c) #1 { 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 e0d0197c6ead5..fb4c653b709f3 100644 --- a/llvm/test/CodeGen/NVPTX/misched_func_call.ll +++ b/llvm/test/CodeGen/NVPTX/misched_func_call.ll @@ -9,7 +9,7 @@ define ptx_kernel void @wombat(i32 %arg, i32 %arg1, i32 %arg2) { ; 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 afabc113541c2..c8daf3b5760f5 100644 --- a/llvm/test/CodeGen/NVPTX/param-add.ll +++ b/llvm/test/CodeGen/NVPTX/param-add.ll @@ -15,7 +15,7 @@ define i32 @test(%struct.1float alignstack(32) %data) { ; 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 31fd8ebae1a81..0b020b7751387 100644 --- a/llvm/test/CodeGen/NVPTX/rcp-opt.ll +++ b/llvm/test/CodeGen/NVPTX/rcp-opt.ll @@ -9,7 +9,7 @@ target triple = "nvptx64-nvidia-cuda" 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 @test1(double %in) { 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 @test2(double %in) { 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 2a12e9b364a54..020a61a1675aa 100644 --- a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll +++ b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll @@ -115,7 +115,7 @@ define half @reduce_fadd_half_reassoc_nonpow2(<7 x half> %in) { 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(<8 x float> %in) { 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(<8 x float> %in) { 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 half @reduce_fmul_half_reassoc_nonpow2(<7 x half> %in) { 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(<8 x float> %in) { 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(<8 x float> %in) { 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 @@ define float @reduce_fmax_float(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmax_float_reassoc(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmax_float_reassoc_nonpow2(<7 x float> %in) { ; ; 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 @@ define float @reduce_fmin_float(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmin_float_reassoc(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmin_float_reassoc_nonpow2(<7 x float> %in) { ; ; 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 @@ define float @reduce_fmaximum_float(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmaximum_float_reassoc(<8 x float> %in) { ; ; 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 @@ define float @reduce_fmaximum_float_reassoc_nonpow2(<7 x float> %in) { ; ; 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 @@ define float @reduce_fminimum_float(<8 x float> %in) { ; ; 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 @@ define float @reduce_fminimum_float_reassoc(<8 x float> %in) { ; ; 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 @@ define float @reduce_fminimum_float_reassoc_nonpow2(<7 x float> %in) { ; ; 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 af113e75fd143..ed785298f5900 100644 --- a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll +++ b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll @@ -7,7 +7,7 @@ define float @redux_sync_fmin(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmin_abs(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmin_NaN(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmin_abs_NaN(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmax(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmax_abs(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmax_NaN(float %src, i32 %mask) { ; 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 @@ define float @redux_sync_fmax_abs_NaN(float %src, i32 %mask) { ; 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 cf2433ad75a97..7b4ebcae8a67c 100644 --- a/llvm/test/CodeGen/NVPTX/reg-types.ll +++ b/llvm/test/CodeGen/NVPTX/reg-types.ll @@ -25,9 +25,9 @@ entry: %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 ab1447607ab65..e8ad68909e286 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_ii(float %val) { 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_ir(float %val) { 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_ii(double %val) { 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_ir(double %val) { 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_iiii() { 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_irrr(float %b, float %c, float %d) { 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_rirr(float %a, float %c, float %d) { 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_rrir(float %a, float %b, float %d) { 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_rrri(float %a, float %b, float %c) { 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_iirr(float %c, float %d) { 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_irir(float %b, float %d) { 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_irri(float %b, float %c) { 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_riir(float %a, float %d) { 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_riri(float %a, float %c) { 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_rrii(float %a, float %b) { 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_iiir(float %d) { 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_iiri(float %c) { 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_irii(float %b) { 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 7a7904a2f0425..3afff3245fbf6 100644 --- a/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll +++ b/llvm/test/CodeGen/NVPTX/surf-read-cuda.ll @@ -14,7 +14,7 @@ define ptx_kernel void @foo(i64 %img, ptr %red, i32 %idx) { ; 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 @@ define ptx_kernel void @bar(ptr %red, i32 %idx) { ; 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 61837bde82ece..4e4e3f3aaec62 100644 --- a/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll +++ b/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll @@ -14,7 +14,7 @@ define ptx_kernel void @foo(i64 %img, ptr %red, i32 %idx) { ; 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 @@ define ptx_kernel void @bar(ptr %red, i32 %idx) { ; 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 @@ define ptx_kernel void @baz(ptr %red, i32 %idx) { ; 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 35db4894c1b49..9da361455a656 100644 --- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll +++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll @@ -13,7 +13,7 @@ define dso_local i32 @variadics1(i32 noundef %first, ...) { ; 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 fa2925af37971..1fc945b364c93 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