[llvm] [AMDGPU] Gate fold_pow(x, ±0.5)->sqrt/rsqrt on nsz+ninf fast-math flags (PR #205592)
Wooseok Lee via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 25 07:13:44 PDT 2026
https://github.com/wooseoklee updated https://github.com/llvm/llvm-project/pull/205592
>From cc1d703c0b5e17ee17f4a7c19523fe795b45afbc Mon Sep 17 00:00:00 2001
From: wooseoklee <wolee at amd.com>
Date: Tue, 23 Jun 2026 14:32:27 -0500
Subject: [PATCH] =?UTF-8?q?[AMDGPU]=20Gate=20fold=5Fpow(x,=20=C2=B10.5)->s?=
=?UTF-8?q?qrt/rsqrt=20on=20nsz+ninf=20fast-math=20flags?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
AMDGPULibCalls::fold_pow rewrote pow(x, 0.5)->sqrt(x) and
pow(x, -0.5)->rsqrt(x) unconditionally, firing before the
isUnsafeFiniteOnlyMath guard. This is incorrect for two IEEE corner
cases:
pow(-Inf, 0.5) == +Inf but sqrt(-Inf) == NaN
pow(-0.0, 0.5) == +0.0 but sqrt(-0.0) == -0.0
Guard the fold on hasNoSignedZeros() && (IsPowr || hasNoInfs()), where
IsPowr is true for powr/powr_fast. The OpenCL spec requires x >= 0 for
powr, so -Inf is undefined behaviour and the ninf check can be skipped;
-0.0 is a valid input for powr since -0.0 >= 0 by IEEE comparison, so
nsz is still required for all variants. afn alone is not sufficient
since it only permits approximate results and says nothing about the
treatment of infinities or signed zeros.
Update the four autogenerated test files to reflect the new semantics:
- Add nsz (and ninf for pow) to existing afn-only ±0.5 call sites so
those tests continue to exercise the fold as intended. powr tests
use nsz only, documenting that ninf is not required.
- Add afn-only (no nsz/ninf) ±0.5 tests in pow-fast.ll and
powr-fast.ll to document that afn alone does not trigger the fold.
---
llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp | 36 +-
.../amdgpu-simplify-libcall-pow-fast.ll | 58 ++-
.../AMDGPU/amdgpu-simplify-libcall-pow.ll | 483 ++++++++----------
.../amdgpu-simplify-libcall-powr-fast.ll | 40 +-
.../AMDGPU/amdgpu-simplify-libcall-powr.ll | 196 +++----
5 files changed, 418 insertions(+), 395 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
index f01487effb8ad..7c0a3fb466b14 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
@@ -935,18 +935,30 @@ bool AMDGPULibCalls::fold_pow(FPMathOperator *FPOp, IRBuilder<> &B,
}
if (CF && (CF->isExactlyValue(0.5) || CF->isExactlyValue(-0.5))) {
- // pow[r](x, [-]0.5) = sqrt(x)
- bool issqrt = CF->isExactlyValue(0.5);
- if (FunctionCallee FPExpr =
- getFunction(M, AMDGPULibFunc(issqrt ? AMDGPULibFunc::EI_SQRT
- : AMDGPULibFunc::EI_RSQRT,
- FInfo))) {
- LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << FInfo.getName()
- << '(' << *opr0 << ")\n");
- Value *nval = CreateCallEx(B,FPExpr, opr0, issqrt ? "__pow2sqrt"
- : "__pow2rsqrt");
- replaceCall(FPOp, nval);
- return true;
+ // pow[r](x, [-]0.5) = sqrt(x) / rsqrt(x)
+ //
+ // sqrt/rsqrt and pow disagree on two negative inputs:
+ // pow(-Inf, 0.5) == +Inf but sqrt(-Inf) == NaN (ninf case)
+ // pow(-0.0, 0.5) == +0.0 but sqrt(-0.0) == -0.0 (nsz case)
+ // powr requires x >= 0 by the OpenCL spec, so -Inf is undefined behaviour
+ // and the ninf check can be skipped for powr/powr_fast. -0.0 is a valid
+ // input for powr since -0.0 >= 0 by IEEE comparison, so nsz is still
+ // required for all variants.
+ bool IsPowr = FInfo.getId() == AMDGPULibFunc::EI_POWR ||
+ FInfo.getId() == AMDGPULibFunc::EI_POWR_FAST;
+ if (FPOp->hasNoSignedZeros() && (IsPowr || FPOp->hasNoInfs())) {
+ bool issqrt = CF->isExactlyValue(0.5);
+ if (FunctionCallee FPExpr =
+ getFunction(M, AMDGPULibFunc(issqrt ? AMDGPULibFunc::EI_SQRT
+ : AMDGPULibFunc::EI_RSQRT,
+ FInfo))) {
+ LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << FInfo.getName()
+ << '(' << *opr0 << ")\n");
+ Value *nval = CreateCallEx(B, FPExpr, opr0,
+ issqrt ? "__pow2sqrt" : "__pow2rsqrt");
+ replaceCall(FPOp, nval);
+ return true;
+ }
}
}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-fast.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-fast.ll
index 961412ae45d2c..87e492cb6bef7 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-fast.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow-fast.ll
@@ -228,11 +228,11 @@ define float @test_pow_afn_f32__half(float %x) #0 {
; CHECK-LABEL: define float @test_pow_afn_f32__half(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call afn float @_Z4sqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn float @_Z4sqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2SQRT]]
;
entry:
- %call = tail call afn float @_Z3powff(float %x, float 0.5)
+ %call = tail call nsz ninf afn float @_Z3powff(float %x, float 0.5)
ret float %call
}
@@ -240,14 +240,40 @@ define float @test_pow_afn_f32__neghalf(float %x) #0 {
; CHECK-LABEL: define float @test_pow_afn_f32__neghalf(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn float @_Z5rsqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn float @_Z5rsqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2RSQRT]]
;
+entry:
+ %call = tail call nsz ninf afn float @_Z3powff(float %x, float -0.5)
+ ret float %call
+}
+
+; pow(x, -0.5) with afn only must NOT fold to rsqrt: nsz and ninf are required.
+define float @test_pow_afn_only_f32__neghalf(float %x) #0 {
+; CHECK-LABEL: define float @test_pow_afn_only_f32__neghalf(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z10__pow_fastff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT: ret float [[CALL]]
+;
entry:
%call = tail call afn float @_Z3powff(float %x, float -0.5)
ret float %call
}
+; pow(x, 0.5) with afn only must NOT fold to sqrt: nsz and ninf are required.
+define float @test_pow_afn_only_f32__half(float %x) #0 {
+; CHECK-LABEL: define float @test_pow_afn_only_f32__half(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z10__pow_fastff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT: ret float [[CALL]]
+;
+entry:
+ %call = tail call afn float @_Z3powff(float %x, float 0.5)
+ ret float %call
+}
+
define <2 x float> @test_pow_afn_v2f32__0(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_pow_afn_v2f32__0(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
@@ -322,11 +348,11 @@ define <2 x float> @test_pow_afn_v2f32__half(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_pow_afn_v2f32__half(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
; CHECK-NEXT: ret <2 x float> [[__POW2SQRT]]
;
entry:
- %call = tail call afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> splat (float 0.5))
+ %call = tail call nsz ninf afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> splat (float 0.5))
ret <2 x float> %call
}
@@ -334,11 +360,11 @@ define <2 x float> @test_pow_afn_v2f32__neghalf(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_pow_afn_v2f32__neghalf(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
; CHECK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
entry:
- %call = tail call afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
+ %call = tail call nsz ninf afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
ret <2 x float> %call
}
@@ -527,7 +553,7 @@ define float @test__pow_fast_f32__half(float %x) #0 {
; CHECK-LABEL: define float @test__pow_fast_f32__half(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call float @_Z4sqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = tail call float @_Z10__pow_fastff(float [[X]], float 5.000000e-01)
; CHECK-NEXT: ret float [[__POW2SQRT]]
;
entry:
@@ -539,11 +565,11 @@ define float @test__pow_fast_afn_f32__half(float %x) #0 {
; CHECK-LABEL: define float @test__pow_fast_afn_f32__half(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call afn float @_Z4sqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn float @_Z4sqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2SQRT]]
;
entry:
- %call = tail call afn float @_Z10__pow_fastff(float %x, float 0.5)
+ %call = tail call nsz ninf afn float @_Z10__pow_fastff(float %x, float 0.5)
ret float %call
}
@@ -551,7 +577,7 @@ define float @test__pow_fast_f32__neghalf(float %x) #0 {
; CHECK-LABEL: define float @test__pow_fast_f32__neghalf(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call float @_Z5rsqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = tail call float @_Z10__pow_fastff(float [[X]], float -5.000000e-01)
; CHECK-NEXT: ret float [[__POW2RSQRT]]
;
entry:
@@ -563,11 +589,11 @@ define float @test__pow_fast_afn_f32__neghalf(float %x) #0 {
; CHECK-LABEL: define float @test__pow_fast_afn_f32__neghalf(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn float @_Z5rsqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn float @_Z5rsqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2RSQRT]]
;
entry:
- %call = tail call afn float @_Z10__pow_fastff(float %x, float -0.5)
+ %call = tail call nsz ninf afn float @_Z10__pow_fastff(float %x, float -0.5)
ret float %call
}
@@ -575,7 +601,7 @@ define <2 x float> @test__pow_fast_v2f32__neghalf(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test__pow_fast_v2f32__neghalf(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = tail call <2 x float> @_Z10__pow_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
; CHECK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
entry:
@@ -587,11 +613,11 @@ define <2 x float> @test__pow_fast_afn_v2f32__neghalf(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test__pow_fast_afn_v2f32__neghalf(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
; CHECK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
entry:
- %call = tail call afn <2 x float> @_Z10__pow_fastDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
+ %call = tail call nsz ninf afn <2 x float> @_Z10__pow_fastDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
ret <2 x float> %call
}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
index 5ad756f69a371..cb7e990a7a62c 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
@@ -1371,309 +1371,288 @@ define <3 x float> @test_pow_afn_v3f32_neg0.0_splat_undef(<3 x float> %x, <3 x f
define float @test_pow_afn_f32_0.5(float %x) {
; PRELINK-LABEL: define float @test_pow_afn_f32_0.5
; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn float @_Z4sqrtf(float [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn float @_Z4sqrtf(float [[X]])
; PRELINK-NEXT: ret float [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define float @test_pow_afn_f32_0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq float [[X]], 1.000000e+00
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn i1 [[TMP1]], float 1.000000e+00, float 5.000000e-01
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn float @llvm.fabs.f32(float [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn float @llvm.log2.f32(float [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn float [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn float @llvm.trunc.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq float [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn float [[TMP2]], 5.000000e-01
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn float @llvm.trunc.f32(float [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une float [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq float [[X]], 1.000000e+00
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn i1 [[TMP1]], float 1.000000e+00, float 5.000000e-01
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn float @llvm.fabs.f32(float [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn float @llvm.log2.f32(float [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn float [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq float [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn float [[TMP2]], 5.000000e-01
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une float [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and i1 [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn i1 [[TMP12]], float [[X]], float 1.000000e+00
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn float @llvm.copysign.f32(float [[TMP6]], float [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn float @llvm.trunc.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une float [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn i1 [[TMP12]], float [[X]], float 1.000000e+00
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn float @llvm.copysign.f32(float [[TMP6]], float [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une float [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt float [[X]], 0.000000e+00
; NOPRELINK-NEXT: [[TMP18:%.*]] = and i1 [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn i1 [[TMP18]], float +qnan, float [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn float @llvm.fabs.f32(float [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq float [[TMP20]], +inf
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or i1 [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = select afn i1 [[TMP22]], float 0.000000e+00, float +inf
-; NOPRELINK-NEXT: [[TMP25:%.*]] = select afn i1 [[TMP12]], float [[X]], float 0.000000e+00
-; NOPRELINK-NEXT: [[TMP26:%.*]] = call afn float @llvm.copysign.f32(float [[TMP24]], float [[TMP25]])
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn i1 [[TMP23]], float [[TMP26]], float [[TMP19]]
-; NOPRELINK-NEXT: [[TMP28:%.*]] = fcmp afn uno float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn i1 [[TMP28]], float +qnan, float [[TMP27]]
-; NOPRELINK-NEXT: ret float [[TMP29]]
-;
- %pow = tail call afn float @_Z3powff(float %x, float 0.5)
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn i1 [[TMP18]], float +qnan, float [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP21:%.*]] = select ninf nsz afn i1 [[TMP20]], float 0.000000e+00, float +inf
+; NOPRELINK-NEXT: [[TMP22:%.*]] = select ninf nsz afn i1 [[TMP12]], float [[X]], float 0.000000e+00
+; NOPRELINK-NEXT: [[TMP23:%.*]] = call ninf nsz afn float @llvm.copysign.f32(float [[TMP21]], float [[TMP22]])
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn i1 [[TMP20]], float [[TMP23]], float [[TMP19]]
+; NOPRELINK-NEXT: [[TMP25:%.*]] = fcmp ninf nsz afn uno float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn i1 [[TMP25]], float +qnan, float [[TMP24]]
+; NOPRELINK-NEXT: ret float [[TMP26]]
+;
+ %pow = tail call nsz ninf afn float @_Z3powff(float %x, float 0.5)
ret float %pow
}
define float @test_pow_afn_f32_neg0.5(float %x) {
; PRELINK-LABEL: define float @test_pow_afn_f32_neg0.5
; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn float @_Z5rsqrtf(float [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn float @_Z5rsqrtf(float [[X]])
; PRELINK-NEXT: ret float [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define float @test_pow_afn_f32_neg0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq float [[X]], 1.000000e+00
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn i1 [[TMP1]], float 1.000000e+00, float -5.000000e-01
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn float @llvm.fabs.f32(float [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn float @llvm.log2.f32(float [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn float [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn float @llvm.trunc.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq float [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn float [[TMP2]], 5.000000e-01
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn float @llvm.trunc.f32(float [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une float [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq float [[X]], 1.000000e+00
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn i1 [[TMP1]], float 1.000000e+00, float -5.000000e-01
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn float @llvm.fabs.f32(float [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn float @llvm.log2.f32(float [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn float [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq float [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn float [[TMP2]], 5.000000e-01
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une float [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and i1 [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn i1 [[TMP12]], float [[X]], float 1.000000e+00
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn float @llvm.copysign.f32(float [[TMP6]], float [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn float @llvm.trunc.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une float [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn i1 [[TMP12]], float [[X]], float 1.000000e+00
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn float @llvm.copysign.f32(float [[TMP6]], float [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn float @llvm.trunc.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une float [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt float [[X]], 0.000000e+00
; NOPRELINK-NEXT: [[TMP18:%.*]] = and i1 [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn i1 [[TMP18]], float +qnan, float [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn float @llvm.fabs.f32(float [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq float [[TMP20]], +inf
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or i1 [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = fcmp afn olt float [[TMP2]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP25:%.*]] = xor i1 [[TMP22]], [[TMP24]]
-; NOPRELINK-NEXT: [[TMP26:%.*]] = select afn i1 [[TMP25]], float 0.000000e+00, float +inf
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn i1 [[TMP12]], float [[X]], float 0.000000e+00
-; NOPRELINK-NEXT: [[TMP28:%.*]] = call afn float @llvm.copysign.f32(float [[TMP26]], float [[TMP27]])
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn i1 [[TMP23]], float [[TMP28]], float [[TMP19]]
-; NOPRELINK-NEXT: [[TMP30:%.*]] = fcmp afn uno float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP31:%.*]] = select afn i1 [[TMP30]], float +qnan, float [[TMP29]]
-; NOPRELINK-NEXT: ret float [[TMP31]]
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn i1 [[TMP18]], float +qnan, float [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp ninf nsz afn olt float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP22:%.*]] = xor i1 [[TMP20]], [[TMP21]]
+; NOPRELINK-NEXT: [[TMP23:%.*]] = select ninf nsz afn i1 [[TMP22]], float 0.000000e+00, float +inf
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn i1 [[TMP12]], float [[X]], float 0.000000e+00
+; NOPRELINK-NEXT: [[TMP25:%.*]] = call ninf nsz afn float @llvm.copysign.f32(float [[TMP23]], float [[TMP24]])
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn i1 [[TMP20]], float [[TMP25]], float [[TMP19]]
+; NOPRELINK-NEXT: [[TMP27:%.*]] = fcmp ninf nsz afn uno float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP28:%.*]] = select ninf nsz afn i1 [[TMP27]], float +qnan, float [[TMP26]]
+; NOPRELINK-NEXT: ret float [[TMP28]]
;
- %pow = tail call afn float @_Z3powff(float %x, float -0.5)
+ %pow = tail call nsz ninf afn float @_Z3powff(float %x, float -0.5)
ret float %pow
}
define <2 x float> @test_pow_afn_v2f32_0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
; PRELINK-NEXT: ret <2 x float> [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <2 x float> [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une <2 x float> [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn <2 x float> [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and <2 x i1> [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une <2 x float> [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt <2 x float> [[X]], zeroinitializer
; NOPRELINK-NEXT: [[TMP18:%.*]] = and <2 x i1> [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq <2 x float> [[TMP20]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or <2 x i1> [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = select afn <2 x i1> [[TMP22]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
-; NOPRELINK-NEXT: [[TMP25:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
-; NOPRELINK-NEXT: [[TMP26:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP24]], <2 x float> [[TMP25]])
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn <2 x i1> [[TMP23]], <2 x float> [[TMP26]], <2 x float> [[TMP19]]
-; NOPRELINK-NEXT: [[TMP28:%.*]] = fcmp afn uno <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn <2 x i1> [[TMP28]], <2 x float> splat (float +qnan), <2 x float> [[TMP27]]
-; NOPRELINK-NEXT: ret <2 x float> [[TMP29]]
-;
- %pow = tail call afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float 0.5>)
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP21:%.*]] = select ninf nsz afn <2 x i1> [[TMP20]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
+; NOPRELINK-NEXT: [[TMP22:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
+; NOPRELINK-NEXT: [[TMP23:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP21]], <2 x float> [[TMP22]])
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn <2 x i1> [[TMP20]], <2 x float> [[TMP23]], <2 x float> [[TMP19]]
+; NOPRELINK-NEXT: [[TMP25:%.*]] = fcmp ninf nsz afn uno <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn <2 x i1> [[TMP25]], <2 x float> splat (float +qnan), <2 x float> [[TMP24]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP26]]
+;
+ %pow = tail call nsz ninf afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float 0.5>)
ret <2 x float> %pow
}
define <2 x float> @test_pow_afn_v2f32_neg0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_neg0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
; PRELINK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_neg0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> splat (float -5.000000e-01)
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <2 x float> [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une <2 x float> [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> splat (float -5.000000e-01)
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn <2 x float> [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and <2 x i1> [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une <2 x float> [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt <2 x float> [[X]], zeroinitializer
; NOPRELINK-NEXT: [[TMP18:%.*]] = and <2 x i1> [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq <2 x float> [[TMP20]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or <2 x i1> [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = fcmp afn olt <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP25:%.*]] = xor <2 x i1> [[TMP22]], [[TMP24]]
-; NOPRELINK-NEXT: [[TMP26:%.*]] = select afn <2 x i1> [[TMP25]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
-; NOPRELINK-NEXT: [[TMP28:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP26]], <2 x float> [[TMP27]])
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn <2 x i1> [[TMP23]], <2 x float> [[TMP28]], <2 x float> [[TMP19]]
-; NOPRELINK-NEXT: [[TMP30:%.*]] = fcmp afn uno <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP31:%.*]] = select afn <2 x i1> [[TMP30]], <2 x float> splat (float +qnan), <2 x float> [[TMP29]]
-; NOPRELINK-NEXT: ret <2 x float> [[TMP31]]
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp ninf nsz afn olt <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP22:%.*]] = xor <2 x i1> [[TMP20]], [[TMP21]]
+; NOPRELINK-NEXT: [[TMP23:%.*]] = select ninf nsz afn <2 x i1> [[TMP22]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
+; NOPRELINK-NEXT: [[TMP25:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP23]], <2 x float> [[TMP24]])
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn <2 x i1> [[TMP20]], <2 x float> [[TMP25]], <2 x float> [[TMP19]]
+; NOPRELINK-NEXT: [[TMP27:%.*]] = fcmp ninf nsz afn uno <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP28:%.*]] = select ninf nsz afn <2 x i1> [[TMP27]], <2 x float> splat (float +qnan), <2 x float> [[TMP26]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP28]]
;
- %pow = tail call afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float -0.5, float -0.5>)
+ %pow = tail call nsz ninf afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float -0.5, float -0.5>)
ret <2 x float> %pow
}
define <2 x float> @test_pow_afn_v2f32_plus_minus_0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_plus_minus_0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[POW:%.*]] = tail call afn <2 x float> @_Z10__pow_fastDv2_fS_(<2 x float> [[X]], <2 x float> <float 5.000000e-01, float -5.000000e-01>)
+; PRELINK-NEXT: [[POW:%.*]] = tail call ninf nsz afn <2 x float> @_Z10__pow_fastDv2_fS_(<2 x float> [[X]], <2 x float> <float 5.000000e-01, float -5.000000e-01>)
; PRELINK-NEXT: ret <2 x float> [[POW]]
;
; NOPRELINK-LABEL: define <2 x float> @test_pow_afn_v2f32_plus_minus_0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> <float 5.000000e-01, float -5.000000e-01>
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <2 x float> [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une <2 x float> [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float 1.000000e+00), <2 x float> <float 5.000000e-01, float -5.000000e-01>
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn <2 x float> [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn <2 x float> [[TMP2]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and <2 x i1> [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une <2 x float> [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP6]], <2 x float> [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn <2 x float> @llvm.trunc.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une <2 x float> [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt <2 x float> [[X]], zeroinitializer
; NOPRELINK-NEXT: [[TMP18:%.*]] = and <2 x i1> [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq <2 x float> [[TMP20]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or <2 x i1> [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = fcmp afn olt <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP25:%.*]] = xor <2 x i1> [[TMP22]], [[TMP24]]
-; NOPRELINK-NEXT: [[TMP26:%.*]] = select afn <2 x i1> [[TMP25]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
-; NOPRELINK-NEXT: [[TMP28:%.*]] = call afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP26]], <2 x float> [[TMP27]])
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn <2 x i1> [[TMP23]], <2 x float> [[TMP28]], <2 x float> [[TMP19]]
-; NOPRELINK-NEXT: [[TMP30:%.*]] = fcmp afn uno <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP31:%.*]] = select afn <2 x i1> [[TMP30]], <2 x float> splat (float +qnan), <2 x float> [[TMP29]]
-; NOPRELINK-NEXT: ret <2 x float> [[TMP31]]
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn <2 x i1> [[TMP18]], <2 x float> splat (float +qnan), <2 x float> [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp ninf nsz afn olt <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP22:%.*]] = xor <2 x i1> [[TMP20]], [[TMP21]]
+; NOPRELINK-NEXT: [[TMP23:%.*]] = select ninf nsz afn <2 x i1> [[TMP22]], <2 x float> zeroinitializer, <2 x float> splat (float +inf)
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn <2 x i1> [[TMP12]], <2 x float> [[X]], <2 x float> zeroinitializer
+; NOPRELINK-NEXT: [[TMP25:%.*]] = call ninf nsz afn <2 x float> @llvm.copysign.v2f32(<2 x float> [[TMP23]], <2 x float> [[TMP24]])
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn <2 x i1> [[TMP20]], <2 x float> [[TMP25]], <2 x float> [[TMP19]]
+; NOPRELINK-NEXT: [[TMP27:%.*]] = fcmp ninf nsz afn uno <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP28:%.*]] = select ninf nsz afn <2 x i1> [[TMP27]], <2 x float> splat (float +qnan), <2 x float> [[TMP26]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP28]]
;
- %pow = tail call afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float -0.5>)
+ %pow = tail call nsz ninf afn <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float -0.5>)
ret <2 x float> %pow
}
define <3 x float> @test_pow_afn_v3f32_0.5_splat_undef(<3 x float> %x, <3 x float> %y) {
; PRELINK-LABEL: define <3 x float> @test_pow_afn_v3f32_0.5_splat_undef
; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn <3 x float> @_Z4sqrtDv3_f(<3 x float> [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call ninf nsz afn <3 x float> @_Z4sqrtDv3_f(<3 x float> [[X]])
; PRELINK-NEXT: ret <3 x float> [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define <3 x float> @test_pow_afn_v3f32_0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq <3 x float> [[X]], splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <3 x i1> [[TMP1]], <3 x float> splat (float 1.000000e+00), <3 x float> <float 5.000000e-01, float poison, float 5.000000e-01>
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <3 x float> [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <3 x float> [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn <3 x float> [[TMP2]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une <3 x float> [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[X]], splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn <3 x i1> [[TMP1]], <3 x float> splat (float 1.000000e+00), <3 x float> <float 5.000000e-01, float poison, float 5.000000e-01>
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn <3 x float> [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn <3 x float> [[TMP2]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une <3 x float> [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and <3 x i1> [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP6]], <3 x float> [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une <3 x float> [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP6]], <3 x float> [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une <3 x float> [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt <3 x float> [[X]], zeroinitializer
; NOPRELINK-NEXT: [[TMP18:%.*]] = and <3 x i1> [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn <3 x i1> [[TMP18]], <3 x float> splat (float +qnan), <3 x float> [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq <3 x float> [[TMP20]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or <3 x i1> [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = select afn <3 x i1> [[TMP22]], <3 x float> zeroinitializer, <3 x float> splat (float +inf)
-; NOPRELINK-NEXT: [[TMP25:%.*]] = select afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> zeroinitializer
-; NOPRELINK-NEXT: [[TMP26:%.*]] = call afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP24]], <3 x float> [[TMP25]])
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn <3 x i1> [[TMP23]], <3 x float> [[TMP26]], <3 x float> [[TMP19]]
-; NOPRELINK-NEXT: [[TMP28:%.*]] = fcmp afn uno <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn <3 x i1> [[TMP28]], <3 x float> splat (float +qnan), <3 x float> [[TMP27]]
-; NOPRELINK-NEXT: ret <3 x float> [[TMP29]]
-;
- %pow = tail call afn <3 x float> @_Z3powDv3_fS_(<3 x float> %x, <3 x float> <float 0.5, float poison, float 0.5>)
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn <3 x i1> [[TMP18]], <3 x float> splat (float +qnan), <3 x float> [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP21:%.*]] = select ninf nsz afn <3 x i1> [[TMP20]], <3 x float> zeroinitializer, <3 x float> splat (float +inf)
+; NOPRELINK-NEXT: [[TMP22:%.*]] = select ninf nsz afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> zeroinitializer
+; NOPRELINK-NEXT: [[TMP23:%.*]] = call ninf nsz afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP21]], <3 x float> [[TMP22]])
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn <3 x i1> [[TMP20]], <3 x float> [[TMP23]], <3 x float> [[TMP19]]
+; NOPRELINK-NEXT: [[TMP25:%.*]] = fcmp ninf nsz afn uno <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn <3 x i1> [[TMP25]], <3 x float> splat (float +qnan), <3 x float> [[TMP24]]
+; NOPRELINK-NEXT: ret <3 x float> [[TMP26]]
+;
+ %pow = tail call nsz ninf afn <3 x float> @_Z3powDv3_fS_(<3 x float> %x, <3 x float> <float 0.5, float poison, float 0.5>)
ret <3 x float> %pow
}
define <3 x float> @test_pow_afn_v3f32_neg0.5_splat_undef(<3 x float> %x, <3 x float> %y) {
; PRELINK-LABEL: define <3 x float> @test_pow_afn_v3f32_neg0.5_splat_undef
; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn <3 x float> @_Z5rsqrtDv3_f(<3 x float> [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call ninf nsz afn <3 x float> @_Z5rsqrtDv3_f(<3 x float> [[X]])
; PRELINK-NEXT: ret <3 x float> [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define <3 x float> @test_pow_afn_v3f32_neg0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn oeq <3 x float> [[X]], splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <3 x i1> [[TMP1]], <3 x float> splat (float 1.000000e+00), <3 x float> <float -5.000000e-01, float poison, float -5.000000e-01>
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <3 x float> [[TMP2]], [[TMP4]]
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <3 x float> [[TMP7]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan afn <3 x float> [[TMP2]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP9]])
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn une <3 x float> [[TMP10]], [[TMP9]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[X]], splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select ninf nsz afn <3 x i1> [[TMP1]], <3 x float> splat (float 1.000000e+00), <3 x float> <float -5.000000e-01, float poison, float -5.000000e-01>
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call ninf nsz afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call ninf nsz afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul ninf nsz afn <3 x float> [[TMP2]], [[TMP4]]
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call ninf nsz afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[TMP7]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fmul nnan ninf nsz afn <3 x float> [[TMP2]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP9]])
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp ninf nsz afn une <3 x float> [[TMP10]], [[TMP9]]
; NOPRELINK-NEXT: [[TMP12:%.*]] = and <3 x i1> [[TMP8]], [[TMP11]]
-; NOPRELINK-NEXT: [[TMP13:%.*]] = select afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> splat (float 1.000000e+00)
-; NOPRELINK-NEXT: [[TMP14:%.*]] = call afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP6]], <3 x float> [[TMP13]])
-; NOPRELINK-NEXT: [[TMP15:%.*]] = call afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp afn une <3 x float> [[TMP15]], [[TMP2]]
-; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp afn olt <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP13:%.*]] = select ninf nsz afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> splat (float 1.000000e+00)
+; NOPRELINK-NEXT: [[TMP14:%.*]] = call ninf nsz afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP6]], <3 x float> [[TMP13]])
+; NOPRELINK-NEXT: [[TMP15:%.*]] = call ninf nsz afn <3 x float> @llvm.trunc.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP16:%.*]] = fcmp ninf nsz afn une <3 x float> [[TMP15]], [[TMP2]]
+; NOPRELINK-NEXT: [[TMP17:%.*]] = fcmp ninf nsz afn olt <3 x float> [[X]], zeroinitializer
; NOPRELINK-NEXT: [[TMP18:%.*]] = and <3 x i1> [[TMP17]], [[TMP16]]
-; NOPRELINK-NEXT: [[TMP19:%.*]] = select afn <3 x i1> [[TMP18]], <3 x float> splat (float +qnan), <3 x float> [[TMP14]]
-; NOPRELINK-NEXT: [[TMP20:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[X]])
-; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp afn oeq <3 x float> [[TMP20]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP22:%.*]] = fcmp afn oeq <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP23:%.*]] = or <3 x i1> [[TMP21]], [[TMP22]]
-; NOPRELINK-NEXT: [[TMP24:%.*]] = fcmp afn olt <3 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP25:%.*]] = xor <3 x i1> [[TMP22]], [[TMP24]]
-; NOPRELINK-NEXT: [[TMP26:%.*]] = select afn <3 x i1> [[TMP25]], <3 x float> zeroinitializer, <3 x float> splat (float +inf)
-; NOPRELINK-NEXT: [[TMP27:%.*]] = select afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> zeroinitializer
-; NOPRELINK-NEXT: [[TMP28:%.*]] = call afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP26]], <3 x float> [[TMP27]])
-; NOPRELINK-NEXT: [[TMP29:%.*]] = select afn <3 x i1> [[TMP23]], <3 x float> [[TMP28]], <3 x float> [[TMP19]]
-; NOPRELINK-NEXT: [[TMP30:%.*]] = fcmp afn uno <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP31:%.*]] = select afn <3 x i1> [[TMP30]], <3 x float> splat (float +qnan), <3 x float> [[TMP29]]
-; NOPRELINK-NEXT: ret <3 x float> [[TMP31]]
-;
- %pow = tail call afn <3 x float> @_Z3powDv3_fS_(<3 x float> %x, <3 x float> <float -0.5, float poison, float -0.5>)
+; NOPRELINK-NEXT: [[TMP19:%.*]] = select ninf nsz afn <3 x i1> [[TMP18]], <3 x float> splat (float +qnan), <3 x float> [[TMP14]]
+; NOPRELINK-NEXT: [[TMP20:%.*]] = fcmp ninf nsz afn oeq <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP21:%.*]] = fcmp ninf nsz afn olt <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP22:%.*]] = xor <3 x i1> [[TMP20]], [[TMP21]]
+; NOPRELINK-NEXT: [[TMP23:%.*]] = select ninf nsz afn <3 x i1> [[TMP22]], <3 x float> zeroinitializer, <3 x float> splat (float +inf)
+; NOPRELINK-NEXT: [[TMP24:%.*]] = select ninf nsz afn <3 x i1> [[TMP12]], <3 x float> [[X]], <3 x float> zeroinitializer
+; NOPRELINK-NEXT: [[TMP25:%.*]] = call ninf nsz afn <3 x float> @llvm.copysign.v3f32(<3 x float> [[TMP23]], <3 x float> [[TMP24]])
+; NOPRELINK-NEXT: [[TMP26:%.*]] = select ninf nsz afn <3 x i1> [[TMP20]], <3 x float> [[TMP25]], <3 x float> [[TMP19]]
+; NOPRELINK-NEXT: [[TMP27:%.*]] = fcmp ninf nsz afn uno <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP28:%.*]] = select ninf nsz afn <3 x i1> [[TMP27]], <3 x float> splat (float +qnan), <3 x float> [[TMP26]]
+; NOPRELINK-NEXT: ret <3 x float> [[TMP28]]
+;
+ %pow = tail call nsz ninf afn <3 x float> @_Z3powDv3_fS_(<3 x float> %x, <3 x float> <float -0.5, float poison, float -0.5>)
ret <3 x float> %pow
}
@@ -3838,30 +3817,20 @@ define float @test_pow_f32__y_n2(float %x) {
}
define float @test_pow_f32__y_half(float %x) {
-; PRELINK-LABEL: define float @test_pow_f32__y_half
-; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call float @_Z4sqrtf(float [[X]])
-; PRELINK-NEXT: ret float [[__POW2SQRT]]
-;
-; NOPRELINK-LABEL: define float @test_pow_f32__y_half
-; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[POW:%.*]] = tail call float @_Z3powff(float [[X]], float 5.000000e-01)
-; NOPRELINK-NEXT: ret float [[POW]]
+; CHECK-LABEL: define float @test_pow_f32__y_half
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call float @_Z3powff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT: ret float [[POW]]
;
%pow = tail call float @_Z3powff(float %x, float 0.5)
ret float %pow
}
define float @test_pow_f32__y_neg_half(float %x) {
-; PRELINK-LABEL: define float @test_pow_f32__y_neg_half
-; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call float @_Z5rsqrtf(float [[X]])
-; PRELINK-NEXT: ret float [[__POW2RSQRT]]
-;
-; NOPRELINK-LABEL: define float @test_pow_f32__y_neg_half
-; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[POW:%.*]] = tail call float @_Z3powff(float [[X]], float -5.000000e-01)
-; NOPRELINK-NEXT: ret float [[POW]]
+; CHECK-LABEL: define float @test_pow_f32__y_neg_half
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call float @_Z3powff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT: ret float [[POW]]
;
%pow = tail call float @_Z3powff(float %x, float -0.5)
ret float %pow
@@ -3980,30 +3949,20 @@ define <2 x float> @test_pow_v2f32__y_n2(<2 x float> %x) {
}
define <2 x float> @test_pow_v2f32__y_half(<2 x float> %x) {
-; PRELINK-LABEL: define <2 x float> @test_pow_v2f32__y_half
-; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
-; PRELINK-NEXT: ret <2 x float> [[__POW2SQRT]]
-;
-; NOPRELINK-LABEL: define <2 x float> @test_pow_v2f32__y_half
-; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[POW:%.*]] = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 5.000000e-01))
-; NOPRELINK-NEXT: ret <2 x float> [[POW]]
+; CHECK-LABEL: define <2 x float> @test_pow_v2f32__y_half
+; CHECK-SAME: (<2 x float> [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 5.000000e-01))
+; CHECK-NEXT: ret <2 x float> [[POW]]
;
%pow = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float 0.5>)
ret <2 x float> %pow
}
define <2 x float> @test_pow_v2f32__y_neg_half(<2 x float> %x) {
-; PRELINK-LABEL: define <2 x float> @test_pow_v2f32__y_neg_half
-; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
-; PRELINK-NEXT: ret <2 x float> [[__POW2RSQRT]]
-;
-; NOPRELINK-LABEL: define <2 x float> @test_pow_v2f32__y_neg_half
-; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[POW:%.*]] = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
-; NOPRELINK-NEXT: ret <2 x float> [[POW]]
+; CHECK-LABEL: define <2 x float> @test_pow_v2f32__y_neg_half
+; CHECK-SAME: (<2 x float> [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
+; CHECK-NEXT: ret <2 x float> [[POW]]
;
%pow = tail call <2 x float> @_Z3powDv2_fS_(<2 x float> %x, <2 x float> <float -0.5, float -0.5>)
ret <2 x float> %pow
@@ -4100,8 +4059,8 @@ define float @test_pow_f32__known_positive__y_2(float nofpclass(ninf nnorm nsub)
define float @test_pow_f32__known_positive__y_half(float nofpclass(ninf nnorm nsub) %x) {
; PRELINK-LABEL: define float @test_pow_f32__known_positive__y_half
; PRELINK-SAME: (float nofpclass(ninf nsub nnorm) [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call float @_Z4sqrtf(float [[X]])
-; PRELINK-NEXT: ret float [[__POW2SQRT]]
+; PRELINK-NEXT: [[POW:%.*]] = tail call float @_Z4powrff(float [[X]], float 5.000000e-01)
+; PRELINK-NEXT: ret float [[POW]]
;
; NOPRELINK-LABEL: define float @test_pow_f32__known_positive__y_half
; NOPRELINK-SAME: (float nofpclass(ninf nsub nnorm) [[X:%.*]]) {
@@ -4115,8 +4074,8 @@ define float @test_pow_f32__known_positive__y_half(float nofpclass(ninf nnorm ns
define float @test_pow_f32__known_positive__y_neghalf(float nofpclass(ninf nnorm nsub) %x) {
; PRELINK-LABEL: define float @test_pow_f32__known_positive__y_neghalf
; PRELINK-SAME: (float nofpclass(ninf nsub nnorm) [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call float @_Z5rsqrtf(float [[X]])
-; PRELINK-NEXT: ret float [[__POW2RSQRT]]
+; PRELINK-NEXT: [[POW:%.*]] = tail call float @_Z4powrff(float [[X]], float -5.000000e-01)
+; PRELINK-NEXT: ret float [[POW]]
;
; NOPRELINK-LABEL: define float @test_pow_f32__known_positive__y_neghalf
; NOPRELINK-SAME: (float nofpclass(ninf nsub nnorm) [[X:%.*]]) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
index 20ad490bdf0f8..489513f7c9cd6 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
@@ -171,11 +171,11 @@ define float @test_powr_afn_f32__half(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__half(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call afn float @_Z4sqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = call nsz afn float @_Z4sqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2SQRT]]
;
entry:
- %call = tail call afn float @_Z4powrff(float %x, float 0.5)
+ %call = tail call nsz afn float @_Z4powrff(float %x, float 0.5)
ret float %call
}
@@ -183,14 +183,40 @@ define float @test_powr_afn_f32__neghalf(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__neghalf(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn float @_Z5rsqrtf(float [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call nsz afn float @_Z5rsqrtf(float [[X]])
; CHECK-NEXT: ret float [[__POW2RSQRT]]
;
+entry:
+ %call = tail call nsz afn float @_Z4powrff(float %x, float -0.5)
+ ret float %call
+}
+
+; powr(x, -0.5) with afn only must NOT fold to rsqrt: nsz is required.
+define float @test_powr_afn_only_f32__neghalf(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_only_f32__neghalf(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT: ret float [[CALL]]
+;
entry:
%call = tail call afn float @_Z4powrff(float %x, float -0.5)
ret float %call
}
+; powr(x, 0.5) with afn only must NOT fold to sqrt: nsz is required.
+define float @test_powr_afn_only_f32__half(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_only_f32__half(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT: ret float [[CALL]]
+;
+entry:
+ %call = tail call afn float @_Z4powrff(float %x, float 0.5)
+ ret float %call
+}
+
define <2 x float> @test_powr_afn_v2f32__0(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__0(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
@@ -265,11 +291,11 @@ define <2 x float> @test_powr_afn_v2f32__half(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__half(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2SQRT:%.*]] = call afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2SQRT:%.*]] = call nsz afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
; CHECK-NEXT: ret <2 x float> [[__POW2SQRT]]
;
entry:
- %call = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float 0.5))
+ %call = tail call nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float 0.5))
ret <2 x float> %call
}
@@ -277,11 +303,11 @@ define <2 x float> @test_powr_afn_v2f32__neghalf(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__neghalf(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; CHECK-NEXT: [[__POW2RSQRT:%.*]] = call nsz afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
; CHECK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
entry:
- %call = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
+ %call = tail call nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
ret <2 x float> %call
}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
index 53ed5e311208f..da5f2df5a035a 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
@@ -816,182 +816,182 @@ define <3 x float> @test_powr_afn_v3f32_neg0.0_splat_undef(<3 x float> %x, <3 x
define float @test_powr_afn_f32_0.5(float %x) {
; PRELINK-LABEL: define float @test_powr_afn_f32_0.5
; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn float @_Z4sqrtf(float [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call nsz afn float @_Z4sqrtf(float [[X]])
; PRELINK-NEXT: ret float [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define float @test_powr_afn_f32_0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn i1 [[TMP1]], float +qnan, float [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn float @llvm.fabs.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn float @llvm.log2.f32(float [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn float [[TMP4]], 5.000000e-01
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq float [[TMP2]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn i1 [[TMP7]], float 0.000000e+00, float [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq float [[TMP2]], +inf
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn i1 [[TMP9]], float +inf, float [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno float [[TMP2]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn i1 [[TMP11]], float +qnan, float [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn i1 [[TMP1]], float +qnan, float [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn float @llvm.fabs.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn float @llvm.log2.f32(float [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn float [[TMP4]], 5.000000e-01
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn i1 [[TMP7]], float 0.000000e+00, float [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq float [[TMP2]], +inf
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn i1 [[TMP9]], float +inf, float [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn i1 [[TMP11]], float +qnan, float [[TMP10]]
; NOPRELINK-NEXT: ret float [[TMP12]]
;
- %powr = tail call afn float @_Z4powrff(float %x, float 0.5)
+ %powr = tail call nsz afn float @_Z4powrff(float %x, float 0.5)
ret float %powr
}
define float @test_powr_afn_f32_neg0.5(float %x) {
; PRELINK-LABEL: define float @test_powr_afn_f32_neg0.5
; PRELINK-SAME: (float [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn float @_Z5rsqrtf(float [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call nsz afn float @_Z5rsqrtf(float [[X]])
; PRELINK-NEXT: ret float [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define float @test_powr_afn_f32_neg0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn i1 [[TMP1]], float +qnan, float [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn float @llvm.fabs.f32(float [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn float @llvm.log2.f32(float [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn float [[TMP4]], -5.000000e-01
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq float [[TMP2]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn i1 [[TMP7]], float +inf, float [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq float [[TMP2]], +inf
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn i1 [[TMP9]], float 0.000000e+00, float [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno float [[TMP2]], 0.000000e+00
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn i1 [[TMP11]], float +qnan, float [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt float [[X]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn i1 [[TMP1]], float +qnan, float [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn float @llvm.fabs.f32(float [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn float @llvm.log2.f32(float [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn float [[TMP4]], -5.000000e-01
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn i1 [[TMP7]], float +inf, float [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq float [[TMP2]], +inf
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn i1 [[TMP9]], float 0.000000e+00, float [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn i1 [[TMP11]], float +qnan, float [[TMP10]]
; NOPRELINK-NEXT: ret float [[TMP12]]
;
- %powr = tail call afn float @_Z4powrff(float %x, float -0.5)
+ %powr = tail call nsz afn float @_Z4powrff(float %x, float -0.5)
ret float %powr
}
define <2 x float> @test_powr_afn_v2f32_0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call nsz afn <2 x float> @_Z4sqrtDv2_f(<2 x float> [[X]])
; PRELINK-NEXT: ret <2 x float> [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP4]], splat (float 5.000000e-01)
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <2 x i1> [[TMP7]], <2 x float> zeroinitializer, <2 x float> [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <2 x i1> [[TMP9]], <2 x float> splat (float +inf), <2 x float> [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn <2 x float> [[TMP4]], splat (float 5.000000e-01)
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn <2 x i1> [[TMP7]], <2 x float> zeroinitializer, <2 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn <2 x i1> [[TMP9]], <2 x float> splat (float +inf), <2 x float> [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
; NOPRELINK-NEXT: ret <2 x float> [[TMP12]]
;
- %powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float 0.5>)
+ %powr = tail call nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float 0.5>)
ret <2 x float> %powr
}
define <2 x float> @test_powr_afn_v2f32_neg0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call nsz afn <2 x float> @_Z5rsqrtDv2_f(<2 x float> [[X]])
; PRELINK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP4]], splat (float -5.000000e-01)
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <2 x i1> [[TMP7]], <2 x float> splat (float +inf), <2 x float> [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <2 x i1> [[TMP9]], <2 x float> zeroinitializer, <2 x float> [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn <2 x float> [[TMP4]], splat (float -5.000000e-01)
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn <2 x i1> [[TMP7]], <2 x float> splat (float +inf), <2 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn <2 x i1> [[TMP9]], <2 x float> zeroinitializer, <2 x float> [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
; NOPRELINK-NEXT: ret <2 x float> [[TMP12]]
;
- %powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float -0.5, float -0.5>)
+ %powr = tail call nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float -0.5, float -0.5>)
ret <2 x float> %powr
}
define <2 x float> @test_powr_afn_v2f32_plus_minus_0.5(<2 x float> %x) {
; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_plus_minus_0.5
; PRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> <float 5.000000e-01, float -5.000000e-01>)
+; PRELINK-NEXT: [[POWR:%.*]] = tail call nsz afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> <float 5.000000e-01, float -5.000000e-01>)
; PRELINK-NEXT: ret <2 x float> [[POWR]]
;
; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_plus_minus_0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt <2 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <2 x float> [[TMP4]], <float 5.000000e-01, float -5.000000e-01>
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <2 x i1> [[TMP7]], <2 x float> <float 0.000000e+00, float +inf>, <2 x float> [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <2 x i1> [[TMP9]], <2 x float> <float +inf, float 0.000000e+00>, <2 x float> [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno <2 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt <2 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn <2 x i1> [[TMP1]], <2 x float> splat (float +qnan), <2 x float> [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn <2 x float> @llvm.fabs.v2f32(<2 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn <2 x float> @llvm.log2.v2f32(<2 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn <2 x float> [[TMP4]], <float 5.000000e-01, float -5.000000e-01>
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn <2 x i1> [[TMP7]], <2 x float> <float 0.000000e+00, float +inf>, <2 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq <2 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn <2 x i1> [[TMP9]], <2 x float> <float +inf, float 0.000000e+00>, <2 x float> [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn <2 x i1> [[TMP11]], <2 x float> splat (float +qnan), <2 x float> [[TMP10]]
; NOPRELINK-NEXT: ret <2 x float> [[TMP12]]
;
- %powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float -0.5>)
+ %powr = tail call nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 0.5, float -0.5>)
ret <2 x float> %powr
}
define <3 x float> @test_powr_afn_v3f32_0.5_splat_undef(<3 x float> %x, <3 x float> %y) {
; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.5_splat_undef
; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call afn <3 x float> @_Z4sqrtDv3_f(<3 x float> [[X]])
+; PRELINK-NEXT: [[__POW2SQRT:%.*]] = call nsz afn <3 x float> @_Z4sqrtDv3_f(<3 x float> [[X]])
; PRELINK-NEXT: ret <3 x float> [[__POW2SQRT]]
;
; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <3 x i1> [[TMP1]], <3 x float> splat (float +qnan), <3 x float> [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <3 x float> [[TMP4]], <float 5.000000e-01, float poison, float 5.000000e-01>
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq <3 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <3 x i1> [[TMP7]], <3 x float> <float 0.000000e+00, float poison, float 0.000000e+00>, <3 x float> [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq <3 x float> [[TMP2]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <3 x i1> [[TMP9]], <3 x float> <float +inf, float poison, float +inf>, <3 x float> [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno <3 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn <3 x i1> [[TMP11]], <3 x float> splat (float +qnan), <3 x float> [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn <3 x i1> [[TMP1]], <3 x float> splat (float +qnan), <3 x float> [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn <3 x float> [[TMP4]], <float 5.000000e-01, float poison, float 5.000000e-01>
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn <3 x i1> [[TMP7]], <3 x float> <float 0.000000e+00, float poison, float 0.000000e+00>, <3 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq <3 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn <3 x i1> [[TMP9]], <3 x float> <float +inf, float poison, float +inf>, <3 x float> [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn <3 x i1> [[TMP11]], <3 x float> splat (float +qnan), <3 x float> [[TMP10]]
; NOPRELINK-NEXT: ret <3 x float> [[TMP12]]
;
- %powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float 0.5, float poison, float 0.5>)
+ %powr = tail call nsz afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float 0.5, float poison, float 0.5>)
ret <3 x float> %powr
}
define <3 x float> @test_powr_afn_v3f32_neg0.5_splat_undef(<3 x float> %x, <3 x float> %y) {
; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.5_splat_undef
; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call afn <3 x float> @_Z5rsqrtDv3_f(<3 x float> [[X]])
+; PRELINK-NEXT: [[__POW2RSQRT:%.*]] = call nsz afn <3 x float> @_Z5rsqrtDv3_f(<3 x float> [[X]])
; PRELINK-NEXT: ret <3 x float> [[__POW2RSQRT]]
;
; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp afn olt <3 x float> [[X]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP2:%.*]] = select afn <3 x i1> [[TMP1]], <3 x float> splat (float +qnan), <3 x float> [[X]]
-; NOPRELINK-NEXT: [[TMP3:%.*]] = call afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[TMP2]])
-; NOPRELINK-NEXT: [[TMP4:%.*]] = call afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
-; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul afn <3 x float> [[TMP4]], <float -5.000000e-01, float poison, float -5.000000e-01>
-; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
-; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn oeq <3 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <3 x i1> [[TMP7]], <3 x float> <float +inf, float poison, float +inf>, <3 x float> [[TMP6]]
-; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn oeq <3 x float> [[TMP2]], splat (float +inf)
-; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <3 x i1> [[TMP9]], <3 x float> <float 0.000000e+00, float poison, float 0.000000e+00>, <3 x float> [[TMP8]]
-; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp afn uno <3 x float> [[TMP2]], zeroinitializer
-; NOPRELINK-NEXT: [[TMP12:%.*]] = select afn <3 x i1> [[TMP11]], <3 x float> splat (float +qnan), <3 x float> [[TMP10]]
+; NOPRELINK-NEXT: [[TMP1:%.*]] = fcmp nsz afn olt <3 x float> [[X]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP2:%.*]] = select nsz afn <3 x i1> [[TMP1]], <3 x float> splat (float +qnan), <3 x float> [[X]]
+; NOPRELINK-NEXT: [[TMP3:%.*]] = call nsz afn <3 x float> @llvm.fabs.v3f32(<3 x float> [[TMP2]])
+; NOPRELINK-NEXT: [[TMP4:%.*]] = call nsz afn <3 x float> @llvm.log2.v3f32(<3 x float> [[TMP3]])
+; NOPRELINK-NEXT: [[TMP5:%.*]] = fmul nsz afn <3 x float> [[TMP4]], <float -5.000000e-01, float poison, float -5.000000e-01>
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call nsz afn <3 x float> @llvm.exp2.v3f32(<3 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp nsz afn oeq <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select nsz afn <3 x i1> [[TMP7]], <3 x float> <float +inf, float poison, float +inf>, <3 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp nsz afn oeq <3 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select nsz afn <3 x i1> [[TMP9]], <3 x float> <float 0.000000e+00, float poison, float 0.000000e+00>, <3 x float> [[TMP8]]
+; NOPRELINK-NEXT: [[TMP11:%.*]] = fcmp nsz afn uno <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP12:%.*]] = select nsz afn <3 x i1> [[TMP11]], <3 x float> splat (float +qnan), <3 x float> [[TMP10]]
; NOPRELINK-NEXT: ret <3 x float> [[TMP12]]
;
- %powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float -0.5, float poison, float -0.5>)
+ %powr = tail call nsz afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float -0.5, float poison, float -0.5>)
ret <3 x float> %powr
}
More information about the llvm-commits
mailing list