[llvm] [AMDGPU] Skip fold_pow constant-exponent shortcuts for powr on possibly-negative base (PR #200579)
Arseniy Obolenskiy via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 8 09:11:39 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/200579
>From eda64863a9bff3ba165da127f5158e2fd30881e6 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Sat, 30 May 2026 16:17:24 +0200
Subject: [PATCH] [AMDGPU] Skip fold_pow constant-exponent shortcuts for powr
on possibly-negative base
OpenCL powr(x, y) returns NaN for x < 0, but the constant-exponent shortcuts in fold_pow folded to finite values
---
llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp | 10 +
.../AMDGPU/amdgpu-simplify-libcall-pow.ll | 42 +++
.../amdgpu-simplify-libcall-powr-fast.ll | 56 ++-
.../AMDGPU/amdgpu-simplify-libcall-powr.ll | 329 ++++++++++++++----
4 files changed, 362 insertions(+), 75 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
index dc4225b94b466..763da387d7def 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
@@ -898,6 +898,16 @@ bool AMDGPULibCalls::fold_pow(FPMathOperator *FPOp, IRBuilder<> &B,
// 0x1111111 means that we don't do anything for this call.
int ci_opr1 = (CINT ? (int)CINT->getSExtValue() : 0x1111111);
+ // OpenCL powr(x<0, y) = NaN, but the folds below would turn it into a
+ // finite number. Skip them unless NaNs are ignored or the base is known
+ // non-negative.
+ if ((CF || CINT) && !FPOp->hasNoNaNs() &&
+ (FInfo.getId() == AMDGPULibFunc::EI_POWR ||
+ FInfo.getId() == AMDGPULibFunc::EI_POWR_FAST) &&
+ !cannotBeOrderedLessThanZero(
+ opr0, SQ.getWithInstruction(cast<Instruction>(FPOp))))
+ return false;
+
if ((CF && CF->isZero()) || (CINT && ci_opr1 == 0)) {
// pow/powr/pown(x, 0) == 1
LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> 1\n");
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
index a553bb10ff29d..564222c0a916f 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
@@ -28,6 +28,7 @@ declare float @llvm.rint.f32(float)
declare float @llvm.nearbyint.f32(float)
declare float @llvm.round.f32(float)
declare float @llvm.roundeven.f32(float)
+declare float @_Z4powrff(float, float)
define float @test_pow_fast_f32(float %x, float %y) {
; PRELINK-LABEL: define float @test_pow_fast_f32
@@ -6150,6 +6151,47 @@ define <2 x float> @test_pow_v2f32_known_integral_constant_vector_poison_elt(<2
ret <2 x float> %pow
}
+define float @test_powr_f32_unknown_base_two(float %x) {
+; CHECK-LABEL: define float @test_powr_f32_unknown_base_two
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call float @_Z4powrff(float [[X]], float 2.000000e+00)
+; CHECK-NEXT: ret float [[POW]]
+;
+ %pow = tail call float @_Z4powrff(float %x, float 2.0)
+ ret float %pow
+}
+
+define float @test_powr_f32_unknown_base_neg_one(float %x) {
+; CHECK-LABEL: define float @test_powr_f32_unknown_base_neg_one
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[POW:%.*]] = tail call float @_Z4powrff(float [[X]], float -1.000000e+00)
+; CHECK-NEXT: ret float [[POW]]
+;
+ %pow = tail call float @_Z4powrff(float %x, float -1.0)
+ ret float %pow
+}
+
+define float @test_powr_f32_fabs_base_two(float %x) {
+; CHECK-LABEL: define float @test_powr_f32_fabs_base_two
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[__POW2:%.*]] = fmul float [[X]], [[X]]
+; CHECK-NEXT: ret float [[__POW2]]
+;
+ %ax = call float @llvm.fabs.f32(float %x)
+ %pow = tail call float @_Z4powrff(float %ax, float 2.0)
+ ret float %pow
+}
+
+define float @test_powr_f32_nnan_base_two(float %x) {
+; CHECK-LABEL: define float @test_powr_f32_nnan_base_two
+; CHECK-SAME: (float [[X:%.*]]) {
+; CHECK-NEXT: [[__POW2:%.*]] = fmul nnan float [[X]], [[X]]
+; CHECK-NEXT: ret float [[__POW2]]
+;
+ %pow = tail call nnan float @_Z4powrff(float %x, float 2.0)
+ ret float %pow
+}
+
attributes #0 = { minsize }
attributes #1 = { noinline }
attributes #2 = { strictfp }
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..a52df02ff0eec 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
@@ -101,7 +101,8 @@ define float @test_powr_afn_f32__0(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__0(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret float 1.000000e+00
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 0.000000e+00)
+; CHECK-NEXT: ret float [[CALL]]
;
entry:
%call = tail call afn float @_Z4powrff(float %x, float 0.0)
@@ -112,7 +113,8 @@ define float @test_powr_afn_f32__1(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__1(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret float [[X]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 1.000000e+00)
+; CHECK-NEXT: ret float [[CALL]]
;
entry:
%call = tail call afn float @_Z4powrff(float %x, float 1.0)
@@ -123,7 +125,7 @@ define float @test_powr_afn_f32__2(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__2(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2:%.*]] = fmul afn float [[X]], [[X]]
+; CHECK-NEXT: [[__POW2:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 2.000000e+00)
; CHECK-NEXT: ret float [[__POW2]]
;
entry:
@@ -159,7 +161,7 @@ define float @test_powr_afn_f32__neg1(float %x) #0 {
; CHECK-LABEL: define float @test_powr_afn_f32__neg1(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POWRECIP:%.*]] = fdiv afn float 1.000000e+00, [[X]]
+; CHECK-NEXT: [[__POWRECIP:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -1.000000e+00)
; CHECK-NEXT: ret float [[__POWRECIP]]
;
entry:
@@ -171,7 +173,7 @@ 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:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
; CHECK-NEXT: ret float [[__POW2SQRT]]
;
entry:
@@ -183,7 +185,7 @@ 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:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
; CHECK-NEXT: ret float [[__POW2RSQRT]]
;
entry:
@@ -195,7 +197,8 @@ 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]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret <2 x float> splat (float 1.000000e+00)
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> zeroinitializer)
+; CHECK-NEXT: ret <2 x float> [[CALL]]
;
entry:
%call = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> zeroinitializer)
@@ -206,7 +209,8 @@ define <2 x float> @test_powr_afn_v2f32__1(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__1(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret <2 x float> [[X]]
+; CHECK-NEXT: [[CALL:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 1.000000e+00))
+; CHECK-NEXT: ret <2 x float> [[CALL]]
;
entry:
%call = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float 1.0))
@@ -217,7 +221,7 @@ define <2 x float> @test_powr_afn_v2f32__2(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__2(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2:%.*]] = fmul afn <2 x float> [[X]], [[X]]
+; CHECK-NEXT: [[__POW2:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 2.000000e+00))
; CHECK-NEXT: ret <2 x float> [[__POW2]]
;
entry:
@@ -253,7 +257,7 @@ define <2 x float> @test_powr_afn_v2f32__neg1(<2 x float> %x) #0 {
; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32__neg1(
; CHECK-SAME: <2 x float> [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POWRECIP:%.*]] = fdiv afn <2 x float> splat (float 1.000000e+00), [[X]]
+; CHECK-NEXT: [[__POWRECIP:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -1.000000e+00))
; CHECK-NEXT: ret <2 x float> [[__POWRECIP]]
;
entry:
@@ -265,7 +269,7 @@ 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:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 5.000000e-01))
; CHECK-NEXT: ret <2 x float> [[__POW2SQRT]]
;
entry:
@@ -277,7 +281,7 @@ 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:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
; CHECK-NEXT: ret <2 x float> [[__POW2RSQRT]]
;
entry:
@@ -385,7 +389,8 @@ define float @test__powr_fast_f32_0(float %x) #0 {
; CHECK-LABEL: define float @test__powr_fast_f32_0(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret float 1.000000e+00
+; CHECK-NEXT: [[CALL:%.*]] = tail call float @_Z11__powr_fastff(float [[X]], float 0.000000e+00)
+; CHECK-NEXT: ret float [[CALL]]
;
entry:
%call = tail call float @_Z11__powr_fastff(float %x, float 0.0)
@@ -396,7 +401,15 @@ define float @test__powr_fast_afn_f32_0(float %x) #0 {
; CHECK-LABEL: define float @test__powr_fast_afn_f32_0(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: ret float 1.000000e+00
+; CHECK-NEXT: [[TMP0:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
+; CHECK-NEXT: [[TMP1:%.*]] = select afn i1 [[TMP0]], float +qnan, float [[X]]
+; CHECK-NEXT: [[TMP2:%.*]] = call afn float @llvm.fabs.f32(float [[TMP1]])
+; CHECK-NEXT: [[TMP3:%.*]] = call afn float @llvm.log2.f32(float [[TMP2]])
+; CHECK-NEXT: [[TMP4:%.*]] = fmul afn float [[TMP3]], 0.000000e+00
+; CHECK-NEXT: [[TMP5:%.*]] = call afn float @llvm.exp2.f32(float [[TMP4]])
+; CHECK-NEXT: [[TMP6:%.*]] = fcmp afn ueq float [[TMP1]], 0.000000e+00
+; CHECK-NEXT: [[TMP7:%.*]] = select afn i1 [[TMP6]], float +qnan, float [[TMP5]]
+; CHECK-NEXT: ret float [[TMP7]]
;
entry:
%call = tail call afn float @_Z11__powr_fastff(float %x, float 0.0)
@@ -407,7 +420,7 @@ define float @test__powr_fast_f32_2(float %x) #0 {
; CHECK-LABEL: define float @test__powr_fast_f32_2(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2:%.*]] = fmul float [[X]], [[X]]
+; CHECK-NEXT: [[__POW2:%.*]] = tail call float @_Z11__powr_fastff(float [[X]], float 2.000000e+00)
; CHECK-NEXT: ret float [[__POW2]]
;
entry:
@@ -419,7 +432,18 @@ define float @test__powr_fast_afn_f32_2(float %x) #0 {
; CHECK-LABEL: define float @test__powr_fast_afn_f32_2(
; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[ENTRY:.*:]]
-; CHECK-NEXT: [[__POW2:%.*]] = fmul afn float [[X]], [[X]]
+; CHECK-NEXT: [[TMP0:%.*]] = fcmp afn olt float [[X]], 0.000000e+00
+; CHECK-NEXT: [[TMP1:%.*]] = select afn i1 [[TMP0]], float +qnan, float [[X]]
+; CHECK-NEXT: [[TMP2:%.*]] = call afn float @llvm.fabs.f32(float [[TMP1]])
+; CHECK-NEXT: [[TMP3:%.*]] = call afn float @llvm.log2.f32(float [[TMP2]])
+; CHECK-NEXT: [[TMP4:%.*]] = fmul afn float [[TMP3]], 2.000000e+00
+; CHECK-NEXT: [[TMP5:%.*]] = call afn float @llvm.exp2.f32(float [[TMP4]])
+; CHECK-NEXT: [[TMP6:%.*]] = fcmp afn oeq float [[TMP1]], 0.000000e+00
+; CHECK-NEXT: [[TMP7:%.*]] = select afn i1 [[TMP6]], float 0.000000e+00, float [[TMP5]]
+; CHECK-NEXT: [[TMP8:%.*]] = fcmp afn oeq float [[TMP1]], +inf
+; CHECK-NEXT: [[TMP9:%.*]] = select afn i1 [[TMP8]], float +inf, float [[TMP7]]
+; CHECK-NEXT: [[TMP10:%.*]] = fcmp afn uno float [[TMP1]], 0.000000e+00
+; CHECK-NEXT: [[__POW2:%.*]] = select afn i1 [[TMP10]], float +qnan, float [[TMP9]]
; CHECK-NEXT: ret float [[__POW2]]
;
entry:
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
index 53ed5e311208f..8ca8f12ec3345 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr.ll
@@ -738,36 +738,88 @@ define float @test_powr_afn_f32_poison(float %x) {
}
define float @test_powr_afn_f32_0.0(float %x) {
-; CHECK-LABEL: define float @test_powr_afn_f32_0.0
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: ret float 1.000000e+00
+; PRELINK-LABEL: define float @test_powr_afn_f32_0.0
+; PRELINK-SAME: (float [[X:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 0.000000e+00)
+; PRELINK-NEXT: ret float [[POWR]]
+;
+; NOPRELINK-LABEL: define float @test_powr_afn_f32_0.0
+; 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]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn ueq float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn i1 [[TMP7]], float +qnan, float [[TMP6]]
+; NOPRELINK-NEXT: ret float [[TMP8]]
;
%powr = tail call afn float @_Z4powrff(float %x, float 0.0)
ret float %powr
}
define float @test_powr_afn_f32_neg0.0(float %x) {
-; CHECK-LABEL: define float @test_powr_afn_f32_neg0.0
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: ret float 1.000000e+00
+; PRELINK-LABEL: define float @test_powr_afn_f32_neg0.0
+; PRELINK-SAME: (float [[X:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -0.000000e+00)
+; PRELINK-NEXT: ret float [[POWR]]
+;
+; NOPRELINK-LABEL: define float @test_powr_afn_f32_neg0.0
+; 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]], -0.000000e+00
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn float @llvm.exp2.f32(float [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn ueq float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn i1 [[TMP7]], float +qnan, float [[TMP6]]
+; NOPRELINK-NEXT: ret float [[TMP8]]
;
%powr = tail call afn float @_Z4powrff(float %x, float -0.0)
ret float %powr
}
define <2 x float> @test_powr_afn_v2f32_0.0(<2 x float> %x) {
-; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.0
-; CHECK-SAME: (<2 x float> [[X:%.*]]) {
-; CHECK-NEXT: ret <2 x float> splat (float 1.000000e+00)
+; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.0
+; 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> zeroinitializer)
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.0
+; 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]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp ueq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <2 x i1> [[TMP7]], <2 x float> splat (float +qnan), <2 x float> [[TMP6]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP8]]
;
%powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 0.0, float 0.0>)
ret <2 x float> %powr
}
define <2 x float> @test_powr_afn_v2f32_neg0.0(<2 x float> %x) {
-; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.0
-; CHECK-SAME: (<2 x float> [[X:%.*]]) {
-; CHECK-NEXT: ret <2 x float> splat (float 1.000000e+00)
+; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.0
+; 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> splat (float -0.000000e+00))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.0
+; 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 -0.000000e+00)
+; NOPRELINK-NEXT: [[TMP6:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP5]])
+; NOPRELINK-NEXT: [[TMP7:%.*]] = fcmp afn ueq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP8:%.*]] = select afn <2 x i1> [[TMP7]], <2 x float> splat (float +qnan), <2 x float> [[TMP6]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP8]]
;
%powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float -0.0, float -0.0>)
ret <2 x float> %powr
@@ -796,18 +848,48 @@ define <2 x float> @test_powr_afn_v2f32_plus_minus_0.0(<2 x float> %x) {
}
define <3 x float> @test_powr_afn_v3f32_0.0_splat_undef(<3 x float> %x, <3 x float> %y) {
-; CHECK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.0_splat_undef
-; CHECK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; CHECK-NEXT: ret <3 x float> splat (float 1.000000e+00)
+; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.0_splat_undef
+; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float 0.000000e+00, float poison, float 0.000000e+00>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.0_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 0.000000e+00, float poison, float 0.000000e+00>
+; 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 +qnan, float poison, float +qnan>, <3 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn uno <3 x float> [[TMP2]], <float 0.000000e+00, float poison, float 0.000000e+00>
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <3 x i1> [[TMP9]], <3 x float> splat (float +qnan), <3 x float> [[TMP8]]
+; NOPRELINK-NEXT: ret <3 x float> [[TMP10]]
;
%powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float 0.0, float poison, float 0.0>)
ret <3 x float> %powr
}
define <3 x float> @test_powr_afn_v3f32_neg0.0_splat_undef(<3 x float> %x, <3 x float> %y) {
-; CHECK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.0_splat_undef
-; CHECK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; CHECK-NEXT: ret <3 x float> splat (float 1.000000e+00)
+; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.0_splat_undef
+; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float -0.000000e+00, float poison, float -0.000000e+00>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.0_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 -0.000000e+00, float poison, float -0.000000e+00>
+; 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 +qnan, float poison, float +qnan>, <3 x float> [[TMP6]]
+; NOPRELINK-NEXT: [[TMP9:%.*]] = fcmp afn uno <3 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP10:%.*]] = select afn <3 x i1> [[TMP9]], <3 x float> splat (float +qnan), <3 x float> [[TMP8]]
+; NOPRELINK-NEXT: ret <3 x float> [[TMP10]]
;
%powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float -0.0, float poison, float -0.0>)
ret <3 x float> %powr
@@ -816,8 +898,8 @@ 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: ret float [[__POW2SQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
+; PRELINK-NEXT: ret float [[POWR]]
;
; NOPRELINK-LABEL: define float @test_powr_afn_f32_0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
@@ -842,8 +924,8 @@ define float @test_powr_afn_f32_0.5(float %x) {
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: ret float [[__POW2RSQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
+; PRELINK-NEXT: ret float [[POWR]]
;
; NOPRELINK-LABEL: define float @test_powr_afn_f32_neg0.5
; NOPRELINK-SAME: (float [[X:%.*]]) {
@@ -868,8 +950,8 @@ define float @test_powr_afn_f32_neg0.5(float %x) {
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: ret <2 x float> [[__POW2SQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 5.000000e-01))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
;
; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
@@ -894,8 +976,8 @@ define <2 x float> @test_powr_afn_v2f32_0.5(<2 x float> %x) {
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: ret <2 x float> [[__POW2RSQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
;
; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg0.5
; NOPRELINK-SAME: (<2 x float> [[X:%.*]]) {
@@ -946,8 +1028,8 @@ define <2 x float> @test_powr_afn_v2f32_plus_minus_0.5(<2 x float> %x) {
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: ret <3 x float> [[__POW2SQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float 5.000000e-01, float poison, float 5.000000e-01>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
;
; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
@@ -972,8 +1054,8 @@ define <3 x float> @test_powr_afn_v3f32_0.5_splat_undef(<3 x float> %x, <3 x flo
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: ret <3 x float> [[__POW2RSQRT]]
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float -5.000000e-01, float poison, float -5.000000e-01>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
;
; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg0.5_splat_undef
; NOPRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
@@ -996,38 +1078,102 @@ define <3 x float> @test_powr_afn_v3f32_neg0.5_splat_undef(<3 x float> %x, <3 x
}
define float @test_powr_afn_f32_1.0(float %x) {
-; CHECK-LABEL: define float @test_powr_afn_f32_1.0
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: ret float [[X]]
+; PRELINK-LABEL: define float @test_powr_afn_f32_1.0
+; PRELINK-SAME: (float [[X:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 1.000000e+00)
+; PRELINK-NEXT: ret float [[POWR]]
+;
+; NOPRELINK-LABEL: define float @test_powr_afn_f32_1.0
+; 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:%.*]] = call afn float @llvm.exp2.f32(float [[TMP4]])
+; NOPRELINK-NEXT: [[TMP6:%.*]] = fcmp afn oeq float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP7:%.*]] = select afn i1 [[TMP6]], float 0.000000e+00, float [[TMP5]]
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq float [[TMP2]], +inf
+; NOPRELINK-NEXT: [[TMP9:%.*]] = select afn i1 [[TMP8]], float +inf, float [[TMP7]]
+; NOPRELINK-NEXT: [[TMP10:%.*]] = fcmp afn uno float [[TMP2]], 0.000000e+00
+; NOPRELINK-NEXT: [[TMP11:%.*]] = select afn i1 [[TMP10]], float +qnan, float [[TMP9]]
+; NOPRELINK-NEXT: ret float [[TMP11]]
;
%powr = tail call afn float @_Z4powrff(float %x, float 1.0)
ret float %powr
}
define float @test_powr_afn_f32_neg1.0(float %x) {
-; CHECK-LABEL: define float @test_powr_afn_f32_neg1.0
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: [[__POWRECIP:%.*]] = fdiv afn float 1.000000e+00, [[X]]
-; CHECK-NEXT: ret float [[__POWRECIP]]
+; PRELINK-LABEL: define float @test_powr_afn_f32_neg1.0
+; PRELINK-SAME: (float [[X:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -1.000000e+00)
+; PRELINK-NEXT: ret float [[POWR]]
+;
+; NOPRELINK-LABEL: define float @test_powr_afn_f32_neg1.0
+; 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:%.*]] = fneg afn float [[TMP4]]
+; 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: ret float [[TMP12]]
;
%powr = tail call afn float @_Z4powrff(float %x, float -1.0)
ret float %powr
}
define <2 x float> @test_powr_afn_v2f32_1.0(<2 x float> %x) {
-; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32_1.0
-; CHECK-SAME: (<2 x float> [[X:%.*]]) {
-; CHECK-NEXT: ret <2 x float> [[X]]
+; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_1.0
+; 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> splat (float 1.000000e+00))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_1.0
+; 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:%.*]] = call afn <2 x float> @llvm.exp2.v2f32(<2 x float> [[TMP4]])
+; NOPRELINK-NEXT: [[TMP6:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP7:%.*]] = select afn <2 x i1> [[TMP6]], <2 x float> zeroinitializer, <2 x float> [[TMP5]]
+; NOPRELINK-NEXT: [[TMP8:%.*]] = fcmp afn oeq <2 x float> [[TMP2]], splat (float +inf)
+; NOPRELINK-NEXT: [[TMP9:%.*]] = select afn <2 x i1> [[TMP8]], <2 x float> splat (float +inf), <2 x float> [[TMP7]]
+; NOPRELINK-NEXT: [[TMP10:%.*]] = fcmp afn uno <2 x float> [[TMP2]], zeroinitializer
+; NOPRELINK-NEXT: [[TMP11:%.*]] = select afn <2 x i1> [[TMP10]], <2 x float> splat (float +qnan), <2 x float> [[TMP9]]
+; NOPRELINK-NEXT: ret <2 x float> [[TMP11]]
;
%powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 1.0, float 1.0>)
ret <2 x float> %powr
}
define <2 x float> @test_powr_afn_v2f32_neg1.0(<2 x float> %x) {
-; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg1.0
-; CHECK-SAME: (<2 x float> [[X:%.*]]) {
-; CHECK-NEXT: [[__POWRECIP:%.*]] = fdiv afn <2 x float> splat (float 1.000000e+00), [[X]]
-; CHECK-NEXT: ret <2 x float> [[__POWRECIP]]
+; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg1.0
+; 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> splat (float -1.000000e+00))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_neg1.0
+; 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:%.*]] = fneg afn <2 x float> [[TMP4]]
+; 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: ret <2 x float> [[TMP12]]
;
%powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float -1.0, float -1.0>)
ret <2 x float> %powr
@@ -1060,29 +1206,78 @@ define <2 x float> @test_powr_afn_v2f32_plus_minus_1.0(<2 x float> %x) {
}
define <3 x float> @test_powr_afn_v3f32_1.0_splat_undef(<3 x float> %x, <3 x float> %y) {
-; CHECK-LABEL: define <3 x float> @test_powr_afn_v3f32_1.0_splat_undef
-; CHECK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; CHECK-NEXT: ret <3 x float> [[X]]
+; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_1.0_splat_undef
+; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float 1.000000e+00, float poison, float 1.000000e+00>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_1.0_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 1.000000e+00, float poison, float 1.000000e+00>
+; 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: ret <3 x float> [[TMP12]]
;
%powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float 1.0, float poison, float 1.0>)
ret <3 x float> %powr
}
define <3 x float> @test_powr_afn_v3f32_neg1.0_splat_undef(<3 x float> %x, <3 x float> %y) {
-; CHECK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg1.0_splat_undef
-; CHECK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
-; CHECK-NEXT: [[__POWRECIP:%.*]] = fdiv afn <3 x float> splat (float 1.000000e+00), [[X]]
-; CHECK-NEXT: ret <3 x float> [[__POWRECIP]]
+; PRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg1.0_splat_undef
+; PRELINK-SAME: (<3 x float> [[X:%.*]], <3 x float> [[Y:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn <3 x float> @_Z11__powr_fastDv3_fS_(<3 x float> [[X]], <3 x float> <float -1.000000e+00, float poison, float -1.000000e+00>)
+; PRELINK-NEXT: ret <3 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <3 x float> @test_powr_afn_v3f32_neg1.0_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 -1.000000e+00, float poison, float -1.000000e+00>
+; 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: ret <3 x float> [[TMP12]]
;
%powr = tail call afn <3 x float> @_Z4powrDv3_fS_(<3 x float> %x, <3 x float> <float -1.0, float poison, float -1.0>)
ret <3 x float> %powr
}
define float @test_powr_afn_f32_2.0(float %x) {
-; CHECK-LABEL: define float @test_powr_afn_f32_2.0
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: [[__POW2:%.*]] = fmul afn float [[X]], [[X]]
-; CHECK-NEXT: ret float [[__POW2]]
+; PRELINK-LABEL: define float @test_powr_afn_f32_2.0
+; PRELINK-SAME: (float [[X:%.*]]) {
+; PRELINK-NEXT: [[POWR:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 2.000000e+00)
+; PRELINK-NEXT: ret float [[POWR]]
+;
+; NOPRELINK-LABEL: define float @test_powr_afn_f32_2.0
+; 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]], 2.000000e+00
+; 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: ret float [[TMP12]]
;
%powr = tail call afn float @_Z4powrff(float %x, float 2.0)
ret float %powr
@@ -1115,10 +1310,26 @@ define float @test_powr_afn_f32_neg2.0(float %x) {
}
define <2 x float> @test_powr_afn_v2f32_2.0(<2 x float> %x) {
-; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32_2.0
-; CHECK-SAME: (<2 x float> [[X:%.*]]) {
-; CHECK-NEXT: [[__POW2:%.*]] = fmul afn <2 x float> [[X]], [[X]]
-; CHECK-NEXT: ret <2 x float> [[__POW2]]
+; PRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_2.0
+; 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> splat (float 2.000000e+00))
+; PRELINK-NEXT: ret <2 x float> [[POWR]]
+;
+; NOPRELINK-LABEL: define <2 x float> @test_powr_afn_v2f32_2.0
+; 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 2.000000e+00)
+; 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: ret <2 x float> [[TMP12]]
;
%powr = tail call afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> <float 2.0, float 2.0>)
ret <2 x float> %powr
More information about the llvm-commits
mailing list