[llvm] [NVPTX] Consistently check fast-math flags when lowering fsqrt (PR #143776)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Jun 11 12:52:10 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Alex MacLean (AlexMaclean)
<details>
<summary>Changes</summary>
Ensure that we check the global, function-level, and instruction-level flags when considering whether to use `sqrt.rn` or `sqrt.approx` to lower either `@<!-- -->llvm.sqrt.f32` or `@<!-- -->llvm.nvvm.sqrt.f`
---
Patch is 46.55 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/143776.diff
8 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp (+2-2)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h (+1-1)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+17-7)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.h (+2-1)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (-3)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+16-19)
- (modified) llvm/test/CodeGen/NVPTX/fast-math.ll (+372-95)
- (modified) llvm/test/CodeGen/NVPTX/sqrt-approx.ll (+285-54)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 32223bf3d601e..121ffc5b9d805 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -71,8 +71,8 @@ NVPTXDAGToDAGISel::getDivF32Level(const SDNode *N) const {
return Subtarget->getTargetLowering()->getDivF32Level(*MF, *N);
}
-bool NVPTXDAGToDAGISel::usePrecSqrtF32() const {
- return Subtarget->getTargetLowering()->usePrecSqrtF32();
+bool NVPTXDAGToDAGISel::usePrecSqrtF32(const SDNode *N) const {
+ return Subtarget->getTargetLowering()->usePrecSqrtF32(*MF, N);
}
bool NVPTXDAGToDAGISel::useF32FTZ() const {
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
index 71a5b7ff8cd30..473f4781a6c38 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
@@ -44,7 +44,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
bool doMulWide;
NVPTX::DivPrecisionLevel getDivF32Level(const SDNode *N) const;
- bool usePrecSqrtF32() const;
+ bool usePrecSqrtF32(const SDNode *N) const;
bool useF32FTZ() const;
bool allowFMA() const;
bool allowUnsafeFPMath() const;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d6a134d9abafd..492f4ab76fdbb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -134,14 +134,23 @@ NVPTXTargetLowering::getDivF32Level(const MachineFunction &MF,
return NVPTX::DivPrecisionLevel::IEEE754;
}
-bool NVPTXTargetLowering::usePrecSqrtF32() const {
- if (UsePrecSqrtF32.getNumOccurrences() > 0) {
- // If nvptx-prec-sqrtf32 is used on the command-line, always honor it
+bool NVPTXTargetLowering::usePrecSqrtF32(const MachineFunction &MF,
+ const SDNode *N) const {
+ // If nvptx-prec-sqrtf32 is used on the command-line, always honor it
+ if (UsePrecSqrtF32.getNumOccurrences() > 0)
return UsePrecSqrtF32;
- } else {
- // Otherwise, use sqrt.approx if fast math is enabled
- return !getTargetMachine().Options.UnsafeFPMath;
+
+ // Otherwise, use sqrt.approx if fast math is enabled
+ if (allowUnsafeFPMath(MF))
+ return false;
+
+ if (N) {
+ const SDNodeFlags Flags = N->getFlags();
+ if (Flags.hasApproximateFuncs())
+ return false;
}
+
+ return true;
}
bool NVPTXTargetLowering::useF32FTZ(const MachineFunction &MF) const {
@@ -1134,7 +1143,8 @@ SDValue NVPTXTargetLowering::getSqrtEstimate(SDValue Operand, SelectionDAG &DAG,
bool &UseOneConst,
bool Reciprocal) const {
if (!(Enabled == ReciprocalEstimate::Enabled ||
- (Enabled == ReciprocalEstimate::Unspecified && !usePrecSqrtF32())))
+ (Enabled == ReciprocalEstimate::Unspecified &&
+ !usePrecSqrtF32(DAG.getMachineFunction()))))
return SDValue();
if (ExtraSteps == ReciprocalEstimate::Unspecified)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 8d71022a1f102..0a54a8fd71f32 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -225,7 +225,8 @@ class NVPTXTargetLowering : public TargetLowering {
// Get whether we should use a precise or approximate 32-bit floating point
// sqrt instruction.
- bool usePrecSqrtF32() const;
+ bool usePrecSqrtF32(const MachineFunction &MF,
+ const SDNode *N = nullptr) const;
// Get whether we should use instructions that flush floating-point denormals
// to sign-preserving zero.
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index b646d39194c7e..ecaca4a011053 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -151,9 +151,6 @@ def doRsqrtOpt : Predicate<"doRsqrtOpt()">;
def doMulWide : Predicate<"doMulWide">;
-def do_SQRTF32_APPROX : Predicate<"!usePrecSqrtF32()">;
-def do_SQRTF32_RN : Predicate<"usePrecSqrtF32()">;
-
def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
def hasDotInstructions : Predicate<"Subtarget->hasDotInstructions()">;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index f918160001ba5..1c38cf0cac0ee 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1525,15 +1525,18 @@ def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64", F64RT, F64RT, int_nvvm_sqrt_rz_
def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64", F64RT, F64RT, int_nvvm_sqrt_rm_d>;
def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64", F64RT, F64RT, int_nvvm_sqrt_rp_d>;
+def fsqrt_approx : PatFrags<(ops node:$a),
+ [(fsqrt node:$a),
+ (int_nvvm_sqrt_f node:$a)], [{
+ return !usePrecSqrtF32(N);
+}]>;
+
// nvvm_sqrt intrinsic
-def : Pat<(int_nvvm_sqrt_f f32:$a),
- (INT_NVVM_SQRT_RN_FTZ_F $a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
-def : Pat<(int_nvvm_sqrt_f f32:$a),
- (INT_NVVM_SQRT_RN_F $a)>, Requires<[do_SQRTF32_RN]>;
-def : Pat<(int_nvvm_sqrt_f f32:$a),
- (INT_NVVM_SQRT_APPROX_FTZ_F $a)>, Requires<[doF32FTZ]>;
-def : Pat<(int_nvvm_sqrt_f f32:$a),
- (INT_NVVM_SQRT_APPROX_F $a)>;
+def : Pat<(int_nvvm_sqrt_f f32:$a), (INT_NVVM_SQRT_RN_FTZ_F $a)>, Requires<[doF32FTZ]>;
+def : Pat<(int_nvvm_sqrt_f f32:$a), (INT_NVVM_SQRT_RN_F $a)>;
+
+def : Pat<(fsqrt_approx f32:$a), (INT_NVVM_SQRT_APPROX_FTZ_F $a)>, Requires<[doF32FTZ]>;
+def : Pat<(fsqrt_approx f32:$a), (INT_NVVM_SQRT_APPROX_F $a)>;
//
// Rsqrt
@@ -1556,20 +1559,14 @@ def: Pat<(fdiv f32imm_1, (int_nvvm_sqrt_approx_f f32:$a)),
def: Pat<(fdiv f32imm_1, (int_nvvm_sqrt_approx_ftz_f f32:$a)),
(INT_NVVM_RSQRT_APPROX_FTZ_F $a)>,
Requires<[doRsqrtOpt]>;
-// same for int_nvvm_sqrt_f when non-precision sqrt is requested
-def: Pat<(fdiv f32imm_1, (int_nvvm_sqrt_f f32:$a)),
- (INT_NVVM_RSQRT_APPROX_F $a)>,
- Requires<[doRsqrtOpt, do_SQRTF32_APPROX, doNoF32FTZ]>;
-def: Pat<(fdiv f32imm_1, (int_nvvm_sqrt_f f32:$a)),
- (INT_NVVM_RSQRT_APPROX_FTZ_F $a)>,
- Requires<[doRsqrtOpt, do_SQRTF32_APPROX, doF32FTZ]>;
-def: Pat<(fdiv f32imm_1, (fsqrt f32:$a)),
+// same for int_nvvm_sqrt_f when non-precision sqrt is requested
+def: Pat<(fdiv f32imm_1, (fsqrt_approx f32:$a)),
(INT_NVVM_RSQRT_APPROX_F $a)>,
- Requires<[doRsqrtOpt, do_SQRTF32_APPROX, doNoF32FTZ]>;
-def: Pat<(fdiv f32imm_1, (fsqrt f32:$a)),
+ Requires<[doRsqrtOpt, doNoF32FTZ]>;
+def: Pat<(fdiv f32imm_1, (fsqrt_approx f32:$a)),
(INT_NVVM_RSQRT_APPROX_FTZ_F $a)>,
- Requires<[doRsqrtOpt, do_SQRTF32_APPROX, doF32FTZ]>;
+ Requires<[doRsqrtOpt, doF32FTZ]>;
//
// Add
//
diff --git a/llvm/test/CodeGen/NVPTX/fast-math.ll b/llvm/test/CodeGen/NVPTX/fast-math.ll
index 4cb6a35e796fb..bc48d242f88fd 100644
--- a/llvm/test/CodeGen/NVPTX/fast-math.ll
+++ b/llvm/test/CodeGen/NVPTX/fast-math.ll
@@ -1,58 +1,114 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
declare float @llvm.sqrt.f32(float)
declare double @llvm.sqrt.f64(double)
-; CHECK-LABEL: sqrt_div(
-; CHECK: sqrt.rn.f32
-; CHECK: div.rn.f32
define float @sqrt_div(float %a, float %b) {
+; CHECK-LABEL: sqrt_div(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_param_0];
+; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_param_1];
+; CHECK-NEXT: div.rn.f32 %r4, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
%t1 = tail call float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
}
-; CHECK-LABEL: sqrt_div_fast(
-; CHECK: sqrt.rn.f32
-; CHECK: div.approx.f32
define float @sqrt_div_fast(float %a, float %b) #0 {
+; CHECK-LABEL: sqrt_div_fast(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_param_0];
+; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_param_1];
+; CHECK-NEXT: div.approx.f32 %r4, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
%t1 = tail call float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
}
-; CHECK-LABEL: sqrt_div_fast_ninf(
-; CHECK: sqrt.approx.f32
-; CHECK: div.approx.f32
define float @sqrt_div_fast_ninf(float %a, float %b) #0 {
+; CHECK-LABEL: sqrt_div_fast_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ninf_param_0];
+; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: abs.f32 %r3, %r1;
+; CHECK-NEXT: setp.lt.f32 %p1, %r3, 0f00800000;
+; CHECK-NEXT: selp.f32 %r4, 0f00000000, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r5, [sqrt_div_fast_ninf_param_1];
+; CHECK-NEXT: div.approx.f32 %r6, %r4, %r5;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT: ret;
%t1 = tail call ninf afn float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
}
-; CHECK-LABEL: sqrt_div_ftz(
-; CHECK: sqrt.rn.ftz.f32
-; CHECK: div.rn.ftz.f32
define float @sqrt_div_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: sqrt_div_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_ftz_param_0];
+; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_ftz_param_1];
+; CHECK-NEXT: div.rn.ftz.f32 %r4, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
%t1 = tail call float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
}
-; CHECK-LABEL: sqrt_div_fast_ftz(
-; CHECK: sqrt.rn.ftz.f32
-; CHECK: div.approx.ftz.f32
define float @sqrt_div_fast_ftz(float %a, float %b) #0 #1 {
+; CHECK-LABEL: sqrt_div_fast_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ftz_param_0];
+; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_ftz_param_1];
+; CHECK-NEXT: div.approx.ftz.f32 %r4, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
%t1 = tail call float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
}
-; CHECK-LABEL: sqrt_div_fast_ftz_ninf(
-; CHECK: sqrt.approx.ftz.f32
-; CHECK: div.approx.ftz.f32
define float @sqrt_div_fast_ftz_ninf(float %a, float %b) #0 #1 {
+; CHECK-LABEL: sqrt_div_fast_ftz_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ftz_ninf_param_0];
+; CHECK-NEXT: setp.eq.ftz.f32 %p1, %r1, 0f00000000;
+; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: selp.f32 %r3, 0f00000000, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r4, [sqrt_div_fast_ftz_ninf_param_1];
+; CHECK-NEXT: div.approx.ftz.f32 %r5, %r3, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NEXT: ret;
%t1 = tail call ninf afn float @llvm.sqrt.f32(float %a)
%t2 = fdiv float %t1, %b
ret float %t2
@@ -61,69 +117,117 @@ define float @sqrt_div_fast_ftz_ninf(float %a, float %b) #0 #1 {
; There are no fast-math or ftz versions of sqrt and div for f64. We use
; reciprocal(rsqrt(x)) for sqrt(x), and emit a vanilla divide.
-; CHECK-LABEL: sqrt_div_fast_ftz_f64(
-; CHECK: sqrt.rn.f64
-; CHECK: div.rn.f64
define double @sqrt_div_fast_ftz_f64(double %a, double %b) #0 #1 {
+; CHECK-LABEL: sqrt_div_fast_ftz_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [sqrt_div_fast_ftz_f64_param_0];
+; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;
+; CHECK-NEXT: ld.param.b64 %rd3, [sqrt_div_fast_ftz_f64_param_1];
+; CHECK-NEXT: div.rn.f64 %rd4, %rd2, %rd3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;
+; CHECK-NEXT: ret;
%t1 = tail call double @llvm.sqrt.f64(double %a)
%t2 = fdiv double %t1, %b
ret double %t2
}
-; CHECK-LABEL: sqrt_div_fast_ftz_f64_ninf(
-; CHECK: rsqrt.approx.f64
-; CHECK: rcp.approx.ftz.f64
-; CHECK: div.rn.f64
define double @sqrt_div_fast_ftz_f64_ninf(double %a, double %b) #0 #1 {
+; CHECK-LABEL: sqrt_div_fast_ftz_f64_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [sqrt_div_fast_ftz_f64_ninf_param_0];
+; CHECK-NEXT: abs.f64 %rd2, %rd1;
+; CHECK-NEXT: setp.lt.f64 %p1, %rd2, 0d0010000000000000;
+; CHECK-NEXT: rsqrt.approx.f64 %rd3, %rd1;
+; CHECK-NEXT: rcp.approx.ftz.f64 %rd4, %rd3;
+; CHECK-NEXT: selp.f64 %rd5, 0d0000000000000000, %rd4, %p1;
+; CHECK-NEXT: ld.param.b64 %rd6, [sqrt_div_fast_ftz_f64_ninf_param_1];
+; CHECK-NEXT: div.rn.f64 %rd7, %rd5, %rd6;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;
+; CHECK-NEXT: ret;
%t1 = tail call ninf afn double @llvm.sqrt.f64(double %a)
%t2 = fdiv double %t1, %b
ret double %t2
}
-; CHECK-LABEL: rsqrt(
-; CHECK-NOT: rsqrt.approx
-; CHECK: sqrt.rn.f32
-; CHECK-NOT: rsqrt.approx
define float @rsqrt(float %a) {
+; CHECK-LABEL: rsqrt(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_param_0];
+; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: rcp.rn.f32 %r3, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
%b = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %b
ret float %ret
}
-; CHECK-LABEL: rsqrt_fast(
-; CHECK-NOT: div.
-; CHECK-NOT: sqrt.
-; CHECK: rsqrt.approx.f32
-; CHECK-NOT: div.
-; CHECK-NOT: sqrt.
define float @rsqrt_fast(float %a) #0 {
+; CHECK-LABEL: rsqrt_fast(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_fast_param_0];
+; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%b = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %b
ret float %ret
}
-; CHECK-LABEL: rsqrt_fast_ftz(
-; CHECK-NOT: div.
-; CHECK-NOT: sqrt.
-; CHECK: rsqrt.approx.ftz.f32
-; CHECK-NOT: div.
-; CHECK-NOT: sqrt.
define float @rsqrt_fast_ftz(float %a) #0 #1 {
+; CHECK-LABEL: rsqrt_fast_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_fast_ftz_param_0];
+; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%b = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %b
ret float %ret
}
-; CHECK-LABEL: fadd
-; CHECK: add.rn.f32
define float @fadd(float %a, float %b) {
+; CHECK-LABEL: fadd(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fadd_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [fadd_param_1];
+; CHECK-NEXT: add.rn.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
%t1 = fadd float %a, %b
ret float %t1
}
-; CHECK-LABEL: fadd_ftz
-; CHECK: add.rn.ftz.f32
define float @fadd_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: fadd_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fadd_ftz_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [fadd_ftz_param_1];
+; CHECK-NEXT: add.rn.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
%t1 = fadd float %a, %b
ret float %t1
}
@@ -131,41 +235,83 @@ define float @fadd_ftz(float %a, float %b) #1 {
declare float @llvm.sin.f32(float)
declare float @llvm.cos.f32(float)
-; CHECK-LABEL: fsin_approx_afn
-; CHECK: sin.approx.f32
define float @fsin_approx_afn(float %a) {
+; CHECK-LABEL: fsin_approx_afn(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fsin_approx_afn_param_0];
+; CHECK-NEXT: sin.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%r = tail call afn float @llvm.sin.f32(float %a)
ret float %r
}
-; CHECK-LABEL: fcos_approx_afn
-; CHECK: cos.approx.f32
define float @fcos_approx_afn(float %a) {
+; CHECK-LABEL: fcos_approx_afn(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fcos_approx_afn_param_0];
+; CHECK-NEXT: cos.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%r = tail call afn float @llvm.cos.f32(float %a)
ret float %r
}
-; CHECK-LABEL: fsin_approx
-; CHECK: sin.approx.f32
define float @fsin_approx(float %a) #0 {
+; CHECK-LABEL: fsin_approx(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fsin_approx_param_0];
+; CHECK-NEXT: sin.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%r = tail call float @llvm.sin.f32(float %a)
ret float %r
}
-; CHECK-LABEL: fcos_approx
-; CHECK: cos.approx.f32
define float @fcos_approx(float %a) #0 {
+; CHECK-LABEL: fcos_approx(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [fcos_approx_param_0];
+; CHECK-NEXT: cos.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%r = tail call float @llvm.cos.f32(float %a)
ret float %r
}
-; CHECK-LABEL: repeated_div_recip_allowed
define float @repeated_div_recip_allowed(i1 %pred, float %a, float %b, float %divisor) {
-; CHECK: rcp.rn.f32
-; CHECK: mul.rn.f32
-; CHECK: mul.rn.f32
-; CHECK: mul.rn.f32
-; CHECK: selp.f32
+; CHECK-LABEL: repeated_div_recip_allowed(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_param_0];
+; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_param_3];
+; CHECK-NEXT: rcp.rn.f32 %r3, %r2;
+; CHECK-NEXT: mul.rn.f32 %r4, %r1, %r3;
+; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_recip_allowed_param_2];
+; CHECK-NEXT: mul.rn.f32 %r6, %r5, %r3;
+; CHECK-NEXT: mul.rn.f32 %r7, %r4, %r6;
+; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r8;
+; CHECK-NEXT: ret;
%x = fdiv arcp float %a, %divisor
%y = fdiv arcp float %b, %divisor
%z = fmul float %x, %y
@@ -173,23 +319,51 @@ define float @repeated_div_recip_allowed(i1 %pred, float %a, float %b, float %di
ret float %w
}
-; CHECK-LABEL: repeated_div_recip_allowed_sel
define float @repeated_div_recip_allowed_sel(i1 %pred, float %a, float %b, float %divisor) {
-; CHECK: selp.f32
-; CHECK: div.rn.f32
+; CHECK-LABEL: repeated_div_recip_allowed_sel(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_sel_param_0];
+; CHECK-NEXT: and.b...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/143776
More information about the llvm-commits
mailing list