[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