[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
Thu Jul 9 04:57:21 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 1/4] [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

>From 50aa99eca1f643e33663b7988e2de35c4897faac Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 9 Jul 2026 08:10:08 +0200
Subject: [PATCH 2/4] Address comments

---
 .../amdgpu-simplify-libcall-powr-fast.ll      | 111 ++++++++++++++++--
 1 file changed, 103 insertions(+), 8 deletions(-)

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 a52df02ff0eec..24f966f80e63e 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
@@ -14,6 +14,7 @@ entry:
 }
 
 declare float @_Z4powrff(float, float) #1
+declare float @llvm.fabs.f32(float)
 
 define <2 x float> @test_powr_afn_v2f32(<2 x float> %x, <2 x float> %y) #0 {
 ; CHECK-LABEL: define <2 x float> @test_powr_afn_v2f32(
@@ -133,6 +134,46 @@ entry:
   ret float %call
 }
 
+; powr(x, 2) on an unknown base must NOT fold to x*x: base may be negative.
+define float @test_powr_afn_f32__2_unknown_base(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__2_unknown_base(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 2.000000e+00)
+; CHECK-NEXT:    ret float [[CALL]]
+;
+entry:
+  %call = tail call afn float @_Z4powrff(float %x, float 2.0)
+  ret float %call
+}
+
+; powr(x, 2) with nnan folds to x*x: nnan lets us assume x >= 0.
+define float @test_powr_afn_f32__2_nnan(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__2_nnan(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[__POW2:%.*]] = fmul nnan afn float [[X]], [[X]]
+; CHECK-NEXT:    ret float [[__POW2]]
+;
+entry:
+  %call = tail call nnan afn float @_Z4powrff(float %x, float 2.0)
+  ret float %call
+}
+
+; powr(fabs(x), 2) folds to x*x: the base is known non-negative.
+define float @test_powr_afn_f32__2_fabs_base(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__2_fabs_base(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[__POW2:%.*]] = fmul afn float [[X]], [[X]]
+; CHECK-NEXT:    ret float [[__POW2]]
+;
+entry:
+  %ax = call afn float @llvm.fabs.f32(float %x)
+  %call = tail call afn float @_Z4powrff(float %ax, float 2.0)
+  ret float %call
+}
+
 define float @test_powr_afn_f32__3(float %x) #0 {
 ; CHECK-LABEL: define float @test_powr_afn_f32__3(
 ; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
@@ -169,27 +210,81 @@ entry:
   ret float %call
 }
 
+; powr(x, 0.5) with nnan+nsz folds to sqrt: nnan lets us assume x >= 0.
 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:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT:    [[__POW2SQRT:%.*]] = call nnan 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 nnan nsz afn float @_Z4powrff(float %x, float 0.5)
   ret float %call
 }
 
+; powr(x, 0.5) without nnan must NOT fold to sqrt: base may be negative.
+define float @test_powr_afn_f32__half_no_nnan(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__half_no_nnan(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[CALL:%.*]] = tail call nsz afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT:    ret float [[CALL]]
+;
+entry:
+  %call = tail call nsz afn float @_Z4powrff(float %x, float 0.5)
+  ret float %call
+}
+
+; powr(x, 0.5) without nsz must NOT fold to sqrt: nsz is required.
+define float @test_powr_afn_f32__half_no_nsz(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__half_no_nsz(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[CALL:%.*]] = tail call nnan afn float @_Z11__powr_fastff(float [[X]], float 5.000000e-01)
+; CHECK-NEXT:    ret float [[CALL]]
+;
+entry:
+  %call = tail call nnan afn float @_Z4powrff(float %x, float 0.5)
+  ret float %call
+}
+
+; powr(x, -0.5) with nnan+nsz folds to rsqrt: nnan lets us assume x >= 0.
 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:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT:    [[__POW2RSQRT:%.*]] = call nnan nsz afn float @_Z5rsqrtf(float [[X]])
 ; CHECK-NEXT:    ret float [[__POW2RSQRT]]
 ;
 entry:
-  %call = tail call afn float @_Z4powrff(float %x, float -0.5)
+  %call = tail call nnan nsz afn float @_Z4powrff(float %x, float -0.5)
+  ret float %call
+}
+
+; powr(x, -0.5) without nnan must NOT fold to rsqrt: base may be negative.
+define float @test_powr_afn_f32__neghalf_no_nnan(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__neghalf_no_nnan(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[CALL:%.*]] = tail call nsz afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT:    ret float [[CALL]]
+;
+entry:
+  %call = tail call nsz afn float @_Z4powrff(float %x, float -0.5)
+  ret float %call
+}
+
+; powr(x, -0.5) without nsz must NOT fold to rsqrt: nsz is required.
+define float @test_powr_afn_f32__neghalf_no_nsz(float %x) #0 {
+; CHECK-LABEL: define float @test_powr_afn_f32__neghalf_no_nsz(
+; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT:  [[ENTRY:.*:]]
+; CHECK-NEXT:    [[CALL:%.*]] = tail call nnan afn float @_Z11__powr_fastff(float [[X]], float -5.000000e-01)
+; CHECK-NEXT:    ret float [[CALL]]
+;
+entry:
+  %call = tail call nnan afn float @_Z4powrff(float %x, float -0.5)
   ret float %call
 }
 
@@ -269,11 +364,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:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float 5.000000e-01))
+; CHECK-NEXT:    [[__POW2SQRT:%.*]] = call nnan 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 nnan nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float 0.5))
   ret <2 x float> %call
 }
 
@@ -281,11 +376,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:%.*]] = tail call afn <2 x float> @_Z11__powr_fastDv2_fS_(<2 x float> [[X]], <2 x float> splat (float -5.000000e-01))
+; CHECK-NEXT:    [[__POW2RSQRT:%.*]] = call nnan 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 nnan nsz afn <2 x float> @_Z4powrDv2_fS_(<2 x float> %x, <2 x float> splat (float -0.5))
   ret <2 x float> %call
 }
 

>From ec4ec133f5a33e9e7cbfc6a48e4860be5ab61d97 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 9 Jul 2026 08:34:13 +0200
Subject: [PATCH 3/4] comment 2

---
 .../AMDGPU/amdgpu-simplify-libcall-pow.ll     | 42 -------------------
 .../amdgpu-simplify-libcall-powr-fast.ll      | 13 ------
 2 files changed, 55 deletions(-)

diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
index 9af82cb1619e2..cb7e990a7a62c 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-pow.ll
@@ -28,7 +28,6 @@ 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
@@ -6110,47 +6109,6 @@ 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 24f966f80e63e..86914b1550544 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-simplify-libcall-powr-fast.ll
@@ -134,19 +134,6 @@ entry:
   ret float %call
 }
 
-; powr(x, 2) on an unknown base must NOT fold to x*x: base may be negative.
-define float @test_powr_afn_f32__2_unknown_base(float %x) #0 {
-; CHECK-LABEL: define float @test_powr_afn_f32__2_unknown_base(
-; CHECK-SAME: float [[X:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT:  [[ENTRY:.*:]]
-; CHECK-NEXT:    [[CALL:%.*]] = tail call afn float @_Z11__powr_fastff(float [[X]], float 2.000000e+00)
-; CHECK-NEXT:    ret float [[CALL]]
-;
-entry:
-  %call = tail call afn float @_Z4powrff(float %x, float 2.0)
-  ret float %call
-}
-
 ; powr(x, 2) with nnan folds to x*x: nnan lets us assume x >= 0.
 define float @test_powr_afn_f32__2_nnan(float %x) #0 {
 ; CHECK-LABEL: define float @test_powr_afn_f32__2_nnan(

>From e3840d8d03c660450284c0e8833f44099243f245 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 9 Jul 2026 13:56:45 +0200
Subject: [PATCH 4/4] Address comment

---
 llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp | 124 +++++++++++-----------
 1 file changed, 62 insertions(+), 62 deletions(-)

diff --git a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
index 7a931b50e9270..64e738385c203 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULibCalls.cpp
@@ -901,73 +901,73 @@ bool AMDGPULibCalls::fold_pow(FPMathOperator *FPOp, IRBuilder<> &B,
   // 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) &&
+  bool IsPowr = FInfo.getId() == AMDGPULibFunc::EI_POWR ||
+                FInfo.getId() == AMDGPULibFunc::EI_POWR_FAST;
+  bool SkipConstantFolds =
+      (CF || CINT) && IsPowr && !FPOp->hasNoNaNs() &&
       !cannotBeOrderedLessThanZero(
-          opr0, SQ.getWithInstruction(cast<Instruction>(FPOp))))
-    return false;
+          opr0, SQ.getWithInstruction(cast<Instruction>(FPOp)));
 
-  if ((CF && CF->isZero()) || (CINT && ci_opr1 == 0)) {
-    //  pow/powr/pown(x, 0) == 1
-    LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> 1\n");
-    Constant *cnval = ConstantFP::get(eltType, 1.0);
-    if (getVecSize(FInfo) > 1) {
-      cnval = ConstantDataVector::getSplat(getVecSize(FInfo), cnval);
+  if (!SkipConstantFolds) {
+    if ((CF && CF->isZero()) || (CINT && ci_opr1 == 0)) {
+      //  pow/powr/pown(x, 0) == 1
+      LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> 1\n");
+      Constant *cnval = ConstantFP::get(eltType, 1.0);
+      if (getVecSize(FInfo) > 1) {
+        cnval = ConstantDataVector::getSplat(getVecSize(FInfo), cnval);
+      }
+      replaceCall(FPOp, cnval);
+      return true;
     }
-    replaceCall(FPOp, cnval);
-    return true;
-  }
-  if ((CF && CF->isOne()) || (CINT && ci_opr1 == 1)) {
-    // pow/powr/pown(x, 1.0) = x
-    LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << *opr0 << "\n");
-    replaceCall(FPOp, opr0);
-    return true;
-  }
-  if ((CF && CF->isExactlyValue(2.0)) || (CINT && ci_opr1 == 2)) {
-    // pow/powr/pown(x, 2.0) = x*x
-    LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << *opr0 << " * "
-                      << *opr0 << "\n");
-    Value *nval = B.CreateFMul(opr0, opr0, "__pow2");
-    replaceCall(FPOp, nval);
-    return true;
-  }
-  if ((CF && CF->isMinusOne()) || (CINT && ci_opr1 == -1)) {
-    // pow/powr/pown(x, -1.0) = 1.0/x
-    LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> 1 / " << *opr0 << "\n");
-    Constant *cnval = ConstantFP::get(eltType, 1.0);
-    if (getVecSize(FInfo) > 1) {
-      cnval = ConstantDataVector::getSplat(getVecSize(FInfo), cnval);
+    if ((CF && CF->isOne()) || (CINT && ci_opr1 == 1)) {
+      // pow/powr/pown(x, 1.0) = x
+      LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << *opr0 << "\n");
+      replaceCall(FPOp, opr0);
+      return true;
+    }
+    if ((CF && CF->isExactlyValue(2.0)) || (CINT && ci_opr1 == 2)) {
+      // pow/powr/pown(x, 2.0) = x*x
+      LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> " << *opr0 << " * "
+                        << *opr0 << "\n");
+      Value *nval = B.CreateFMul(opr0, opr0, "__pow2");
+      replaceCall(FPOp, nval);
+      return true;
+    }
+    if ((CF && CF->isMinusOne()) || (CINT && ci_opr1 == -1)) {
+      // pow/powr/pown(x, -1.0) = 1.0/x
+      LLVM_DEBUG(errs() << "AMDIC: " << *FPOp << " ---> 1 / " << *opr0 << "\n");
+      Constant *cnval = ConstantFP::get(eltType, 1.0);
+      if (getVecSize(FInfo) > 1) {
+        cnval = ConstantDataVector::getSplat(getVecSize(FInfo), cnval);
+      }
+      Value *nval = B.CreateFDiv(cnval, opr0, "__powrecip");
+      replaceCall(FPOp, nval);
+      return true;
     }
-    Value *nval = B.CreateFDiv(cnval, opr0, "__powrecip");
-    replaceCall(FPOp, nval);
-    return true;
-  }
 
-  if (CF && (CF->isExactlyValue(0.5) || CF->isExactlyValue(-0.5))) {
-    // 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;
+    if (CF && (CF->isExactlyValue(0.5) || CF->isExactlyValue(-0.5))) {
+      // 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.
+      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;
+        }
       }
     }
   }



More information about the llvm-commits mailing list