[llvm] [NVPTX] Consistently check fast-math flags when lowering fsqrt (PR #143776)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Wed Jun 11 12:51:39 PDT 2025
https://github.com/AlexMaclean created https://github.com/llvm/llvm-project/pull/143776
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`
>From f2c2eba5286f633132d72f8ff2d5ca8789da7e4a Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 11 Jun 2025 15:00:29 +0000
Subject: [PATCH 1/2] pre-commit tests
---
llvm/test/CodeGen/NVPTX/fast-math.ll | 467 ++++++++++++++++++++-----
llvm/test/CodeGen/NVPTX/sqrt-approx.ll | 339 +++++++++++++++---
2 files changed, 657 insertions(+), 149 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/fast-math.ll b/llvm/test/CodeGen/NVPTX/fast-math.ll
index 4cb6a35e796fb..4911e320a73ec 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.rn.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.rn.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.b16 %rs2, %rs1, 1;
+; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_sel_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_sel_param_2];
+; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_recip_allowed_sel_param_3];
+; CHECK-NEXT: div.rn.f32 %r5, %r3, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NEXT: ret;
%x = fdiv arcp float %a, %divisor
%y = fdiv arcp float %b, %divisor
%w = select i1 %pred, float %x, float %y
ret float %w
}
-; CHECK-LABEL: repeated_div_recip_allowed_ftz
define float @repeated_div_recip_allowed_ftz(i1 %pred, float %a, float %b, float %divisor) #1 {
-; CHECK: rcp.rn.ftz.f32
-; CHECK: mul.rn.ftz.f32
-; CHECK: mul.rn.ftz.f32
-; CHECK: mul.rn.ftz.f32
-; CHECK: selp.f32
+; CHECK-LABEL: repeated_div_recip_allowed_ftz(
+; 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_ftz_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_ftz_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_param_3];
+; CHECK-NEXT: rcp.rn.ftz.f32 %r3, %r2;
+; CHECK-NEXT: mul.rn.ftz.f32 %r4, %r1, %r3;
+; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_recip_allowed_ftz_param_2];
+; CHECK-NEXT: mul.rn.ftz.f32 %r6, %r5, %r3;
+; CHECK-NEXT: mul.rn.ftz.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
@@ -197,23 +371,51 @@ define float @repeated_div_recip_allowed_ftz(i1 %pred, float %a, float %b, float
ret float %w
}
-; CHECK-LABEL: repeated_div_recip_allowed_ftz_sel
define float @repeated_div_recip_allowed_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #1 {
-; CHECK: selp.f32
-; CHECK: div.rn.ftz.f32
+; CHECK-LABEL: repeated_div_recip_allowed_ftz_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_ftz_sel_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_ftz_sel_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_sel_param_2];
+; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_recip_allowed_ftz_sel_param_3];
+; CHECK-NEXT: div.rn.ftz.f32 %r5, %r3, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NEXT: ret;
%x = fdiv arcp float %a, %divisor
%y = fdiv arcp float %b, %divisor
%w = select i1 %pred, float %x, float %y
ret float %w
}
-; CHECK-LABEL: repeated_div_fast
define float @repeated_div_fast(i1 %pred, float %a, float %b, float %divisor) #0 {
-; CHECK: rcp.approx.f32
-; CHECK: mul.f32
-; CHECK: mul.f32
-; CHECK: mul.f32
-; CHECK: selp.f32
+; CHECK-LABEL: repeated_div_fast(
+; 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_fast_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_fast_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_param_3];
+; CHECK-NEXT: rcp.approx.f32 %r3, %r2;
+; CHECK-NEXT: mul.f32 %r4, %r1, %r3;
+; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_fast_param_2];
+; CHECK-NEXT: mul.f32 %r6, %r5, %r3;
+; CHECK-NEXT: mul.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 float %a, %divisor
%y = fdiv float %b, %divisor
%z = fmul float %x, %y
@@ -221,23 +423,51 @@ define float @repeated_div_fast(i1 %pred, float %a, float %b, float %divisor) #0
ret float %w
}
-; CHECK-LABEL: repeated_div_fast_sel
define float @repeated_div_fast_sel(i1 %pred, float %a, float %b, float %divisor) #0 {
-; CHECK: selp.f32
-; CHECK: div.approx.f32
+; CHECK-LABEL: repeated_div_fast_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_fast_sel_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_fast_sel_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_sel_param_2];
+; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_fast_sel_param_3];
+; CHECK-NEXT: div.approx.f32 %r5, %r3, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NEXT: ret;
%x = fdiv float %a, %divisor
%y = fdiv float %b, %divisor
%w = select i1 %pred, float %x, float %y
ret float %w
}
-; CHECK-LABEL: repeated_div_fast_ftz
define float @repeated_div_fast_ftz(i1 %pred, float %a, float %b, float %divisor) #0 #1 {
-; CHECK: rcp.approx.ftz.f32
-; CHECK: mul.ftz.f32
-; CHECK: mul.ftz.f32
-; CHECK: mul.ftz.f32
-; CHECK: selp.f32
+; CHECK-LABEL: repeated_div_fast_ftz(
+; 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_fast_ftz_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_fast_ftz_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_param_3];
+; CHECK-NEXT: rcp.approx.ftz.f32 %r3, %r2;
+; CHECK-NEXT: mul.ftz.f32 %r4, %r1, %r3;
+; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_fast_ftz_param_2];
+; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r3;
+; CHECK-NEXT: mul.ftz.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 float %a, %divisor
%y = fdiv float %b, %divisor
%z = fmul float %x, %y
@@ -245,33 +475,80 @@ define float @repeated_div_fast_ftz(i1 %pred, float %a, float %b, float %divisor
ret float %w
}
-; CHECK-LABEL: repeated_div_fast_ftz_sel
define float @repeated_div_fast_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #0 #1 {
-; CHECK: selp.f32
-; CHECK: div.approx.ftz.f32
+; CHECK-LABEL: repeated_div_fast_ftz_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_fast_ftz_sel_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_fast_ftz_sel_param_1];
+; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_sel_param_2];
+; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1;
+; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_fast_ftz_sel_param_3];
+; CHECK-NEXT: div.approx.ftz.f32 %r5, %r3, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NEXT: ret;
%x = fdiv float %a, %divisor
%y = fdiv float %b, %divisor
%w = select i1 %pred, float %x, float %y
ret float %w
}
-; CHECK-LABEL: frem
define float @frem(float %a, float %b) #0 {
- ; CHECK-NOT: testp.infinite
+; CHECK-LABEL: frem(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [frem_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [frem_param_1];
+; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;
+; CHECK-NEXT: neg.f32 %r5, %r4;
+; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT: ret;
%rem = frem float %a, %b
ret float %rem
}
-; CHECK-LABEL: frem_ftz
define float @frem_ftz(float %a, float %b) #0 #1 {
- ; CHECK-NOT: testp.infinite
+; CHECK-LABEL: frem_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [frem_ftz_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [frem_ftz_param_1];
+; CHECK-NEXT: div.approx.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r4, %r3;
+; CHECK-NEXT: neg.ftz.f32 %r5, %r4;
+; CHECK-NEXT: fma.rn.ftz.f32 %r6, %r5, %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT: ret;
%rem = frem float %a, %b
ret float %rem
}
-; CHECK-LABEL: frem_f64
define double @frem_f64(double %a, double %b) #0 {
- ; CHECK-NOT: testp.infinite
+; CHECK-LABEL: frem_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_param_1];
+; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;
+; CHECK-NEXT: neg.f64 %rd5, %rd4;
+; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;
+; CHECK-NEXT: ret;
%rem = frem double %a, %b
ret double %rem
}
diff --git a/llvm/test/CodeGen/NVPTX/sqrt-approx.ll b/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
index e2a3f1cf0d2de..a60ab717f3755 100644
--- a/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
+++ b/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
@@ -1,3 +1,4 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \
; RUN: | FileCheck %s
; RUN: %if ptxas %{ \
@@ -12,34 +13,62 @@ declare double @llvm.sqrt.f64(double)
; -- reciprocal sqrt --
-; CHECK-LABEL: test_rsqrt32
define float @test_rsqrt32(float %a) #0 {
-; CHECK: rsqrt.approx.f32
+; CHECK-LABEL: test_rsqrt32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_param_0];
+; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %val
ret float %ret
}
-; CHECK-LABEL: test_rsqrt_ftz
define float @test_rsqrt_ftz(float %a) #0 #1 {
-; CHECK: rsqrt.approx.ftz.f32
+; CHECK-LABEL: test_rsqrt_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt_ftz_param_0];
+; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %val
ret float %ret
}
-; CHECK-LABEL: test_rsqrt64
define double @test_rsqrt64(double %a) #0 {
-; CHECK: rsqrt.approx.f64
+; CHECK-LABEL: test_rsqrt64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.sqrt.f64(double %a)
%ret = fdiv double 1.0, %val
ret double %ret
}
-; CHECK-LABEL: test_rsqrt64_ftz
-define double @test_rsqrt64_ftz(double %a) #0 #1 {
; There's no rsqrt.approx.ftz.f64 instruction; we just use the non-ftz version.
-; CHECK: rsqrt.approx.f64
+define double @test_rsqrt64_ftz(double %a) #0 #1 {
+; CHECK-LABEL: test_rsqrt64_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_ftz_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.sqrt.f64(double %a)
%ret = fdiv double 1.0, %val
ret double %ret
@@ -47,64 +76,135 @@ define double @test_rsqrt64_ftz(double %a) #0 #1 {
; -- sqrt --
-; CHECK-LABEL: test_sqrt32
define float @test_sqrt32(float %a) #0 {
-; CHECK: sqrt.rn.f32
+; CHECK-LABEL: test_sqrt32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_param_0];
+; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt32_ninf
define float @test_sqrt32_ninf(float %a) #0 {
-; CHECK: sqrt.approx.f32
+; CHECK-LABEL: test_sqrt32_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_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: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt_ftz
define float @test_sqrt_ftz(float %a) #0 #1 {
-; CHECK: sqrt.rn.ftz.f32
+; CHECK-LABEL: test_sqrt_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_ftz_param_0];
+; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt_ftz_ninf
define float @test_sqrt_ftz_ninf(float %a) #0 #1 {
-; CHECK: sqrt.approx.ftz.f32
+; CHECK-LABEL: test_sqrt_ftz_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_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: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt64
define double @test_sqrt64(double %a) #0 {
-; CHECK: sqrt.rn.f64
+; CHECK-LABEL: test_sqrt64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_param_0];
+; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%ret = tail call double @llvm.sqrt.f64(double %a)
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_ninf
-define double @test_sqrt64_ninf(double %a) #0 {
; There's no sqrt.approx.f64 instruction; we emit
; reciprocal(rsqrt.approx.f64(x)). There's no non-ftz approximate reciprocal,
; so we just use the ftz version.
-; CHECK: rsqrt.approx.f64
-; CHECK: rcp.approx.ftz.f64
+define double @test_sqrt64_ninf(double %a) #0 {
+; CHECK-LABEL: test_sqrt64_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_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: st.param.b64 [func_retval0], %rd5;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn double @llvm.sqrt.f64(double %a)
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_ftz
define double @test_sqrt64_ftz(double %a) #0 #1 {
-; CHECK: sqrt.rn.f64
+; CHECK-LABEL: test_sqrt64_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_ftz_param_0];
+; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%ret = tail call double @llvm.sqrt.f64(double %a)
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_ftz_ninf
-define double @test_sqrt64_ftz_ninf(double %a) #0 #1 {
; There's no sqrt.approx.ftz.f64 instruction; we just use the non-ftz version.
-; CHECK: rsqrt.approx.f64
-; CHECK: rcp.approx.ftz.f64
+define double @test_sqrt64_ftz_ninf(double %a) #0 #1 {
+; CHECK-LABEL: test_sqrt64_ftz_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_ftz_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: st.param.b64 [func_retval0], %rd5;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn double @llvm.sqrt.f64(double %a)
ret double %ret
}
@@ -114,93 +214,224 @@ define double @test_sqrt64_ftz_ninf(double %a) #0 #1 {
; The sqrt and rsqrt refinement algorithms both emit an rsqrt.approx, followed
; by some math.
-; CHECK-LABEL: test_rsqrt32_refined
define float @test_rsqrt32_refined(float %a) #0 #2 {
-; CHECK: rsqrt.approx.f32
+; CHECK-LABEL: test_rsqrt32_refined(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_refined_param_0];
+; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: mul.f32 %r3, %r1, %r2;
+; CHECK-NEXT: fma.rn.f32 %r4, %r3, %r2, 0fC0400000;
+; CHECK-NEXT: mul.f32 %r5, %r2, 0fBF000000;
+; CHECK-NEXT: mul.f32 %r6, %r5, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %val
ret float %ret
}
-; CHECK-LABEL: test_sqrt32_refined
define float @test_sqrt32_refined(float %a) #0 #2 {
-; CHECK: sqrt.rn.f32
+; CHECK-LABEL: test_sqrt32_refined(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_param_0];
+; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt32_refined_ninf
define float @test_sqrt32_refined_ninf(float %a) #0 #2 {
-; CHECK: rsqrt.approx.f32
+; CHECK-LABEL: test_sqrt32_refined_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ninf_param_0];
+; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;
+; CHECK-NEXT: mul.f32 %r3, %r1, %r2;
+; CHECK-NEXT: fma.rn.f32 %r4, %r3, %r2, 0fC0400000;
+; CHECK-NEXT: mul.f32 %r5, %r3, 0fBF000000;
+; CHECK-NEXT: mul.f32 %r6, %r5, %r4;
+; CHECK-NEXT: abs.f32 %r7, %r1;
+; CHECK-NEXT: setp.lt.f32 %p1, %r7, 0f00800000;
+; CHECK-NEXT: selp.f32 %r8, 0f00000000, %r6, %p1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r8;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_rsqrt64_refined
define double @test_rsqrt64_refined(double %a) #0 #2 {
-; CHECK: rsqrt.approx.f64
+; CHECK-LABEL: test_rsqrt64_refined(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_refined_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;
+; CHECK-NEXT: mul.f64 %rd5, %rd2, 0dBFE0000000000000;
+; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.sqrt.f64(double %a)
%ret = fdiv double 1.0, %val
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_refined
define double @test_sqrt64_refined(double %a) #0 #2 {
-; CHECK: sqrt.rn.f64
+; CHECK-LABEL: test_sqrt64_refined(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_param_0];
+; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%ret = tail call double @llvm.sqrt.f64(double %a)
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_refined_ninf
define double @test_sqrt64_refined_ninf(double %a) #0 #2 {
-; CHECK: rsqrt.approx.f64
+; CHECK-LABEL: test_sqrt64_refined_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ninf_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;
+; CHECK-NEXT: mul.f64 %rd5, %rd3, 0dBFE0000000000000;
+; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;
+; CHECK-NEXT: abs.f64 %rd7, %rd1;
+; CHECK-NEXT: setp.lt.f64 %p1, %rd7, 0d0010000000000000;
+; CHECK-NEXT: selp.f64 %rd8, 0d0000000000000000, %rd6, %p1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd8;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn double @llvm.sqrt.f64(double %a)
ret double %ret
}
; -- refined sqrt and rsqrt with ftz enabled --
-; CHECK-LABEL: test_rsqrt32_refined_ftz
define float @test_rsqrt32_refined_ftz(float %a) #0 #1 #2 {
-; CHECK: rsqrt.approx.ftz.f32
+; CHECK-LABEL: test_rsqrt32_refined_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_refined_ftz_param_0];
+; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: mul.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: fma.rn.ftz.f32 %r4, %r3, %r2, 0fC0400000;
+; CHECK-NEXT: mul.ftz.f32 %r5, %r2, 0fBF000000;
+; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.sqrt.f32(float %a)
%ret = fdiv float 1.0, %val
ret float %ret
}
-; CHECK-LABEL: test_sqrt32_refined_ftz
define float @test_sqrt32_refined_ftz(float %a) #0 #1 #2 {
-; CHECK: sqrt.rn.ftz.f32
+; CHECK-LABEL: test_sqrt32_refined_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ftz_param_0];
+; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_sqrt32_refined_ftz_ninf
define float @test_sqrt32_refined_ftz_ninf(float %a) #0 #1 #2 {
-; CHECK: rsqrt.approx.ftz.f32
+; CHECK-LABEL: test_sqrt32_refined_ftz_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ftz_ninf_param_0];
+; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;
+; CHECK-NEXT: mul.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: fma.rn.ftz.f32 %r4, %r3, %r2, 0fC0400000;
+; CHECK-NEXT: mul.ftz.f32 %r5, %r3, 0fBF000000;
+; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r4;
+; CHECK-NEXT: setp.eq.ftz.f32 %p1, %r1, 0f00000000;
+; CHECK-NEXT: selp.f32 %r7, 0f00000000, %r6, %p1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn float @llvm.sqrt.f32(float %a)
ret float %ret
}
-; CHECK-LABEL: test_rsqrt64_refined_ftz
-define double @test_rsqrt64_refined_ftz(double %a) #0 #1 #2 {
; There's no rsqrt.approx.ftz.f64, so we just use the non-ftz version.
-; CHECK: rsqrt.approx.f64
+define double @test_rsqrt64_refined_ftz(double %a) #0 #1 #2 {
+; CHECK-LABEL: test_rsqrt64_refined_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_refined_ftz_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;
+; CHECK-NEXT: mul.f64 %rd5, %rd2, 0dBFE0000000000000;
+; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.sqrt.f64(double %a)
%ret = fdiv double 1.0, %val
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_refined_ftz
define double @test_sqrt64_refined_ftz(double %a) #0 #1 #2 {
-; CHECK: sqrt.rn.f64
+; CHECK-LABEL: test_sqrt64_refined_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ftz_param_0];
+; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%ret = tail call double @llvm.sqrt.f64(double %a)
ret double %ret
}
-; CHECK-LABEL: test_sqrt64_refined_ftz_ninf
define double @test_sqrt64_refined_ftz_ninf(double %a) #0 #1 #2 {
-; CHECK: rsqrt.approx.f64
+; CHECK-LABEL: test_sqrt64_refined_ftz_ninf(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ftz_ninf_param_0];
+; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;
+; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;
+; CHECK-NEXT: mul.f64 %rd5, %rd3, 0dBFE0000000000000;
+; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;
+; CHECK-NEXT: abs.f64 %rd7, %rd1;
+; CHECK-NEXT: setp.lt.f64 %p1, %rd7, 0d0010000000000000;
+; CHECK-NEXT: selp.f64 %rd8, 0d0000000000000000, %rd6, %p1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd8;
+; CHECK-NEXT: ret;
%ret = tail call ninf afn double @llvm.sqrt.f64(double %a)
ret double %ret
}
>From 815fc4468fa54bf1825120526d9e7385adf0520b Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 11 Jun 2025 15:15:10 +0000
Subject: [PATCH 2/2] [NVPTX] Consistently check fast-math flags when lowering
fsqrt
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 4 +--
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h | 2 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 24 +++++++++-----
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 3 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 3 --
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 35 ++++++++++-----------
llvm/test/CodeGen/NVPTX/fast-math.ll | 4 +--
llvm/test/CodeGen/NVPTX/sqrt-approx.ll | 8 ++---
8 files changed, 44 insertions(+), 39 deletions(-)
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 4911e320a73ec..bc48d242f88fd 100644
--- a/llvm/test/CodeGen/NVPTX/fast-math.ll
+++ b/llvm/test/CodeGen/NVPTX/fast-math.ll
@@ -29,7 +29,7 @@ define float @sqrt_div_fast(float %a, float %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_param_0];
-; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; 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;
@@ -84,7 +84,7 @@ define float @sqrt_div_fast_ftz(float %a, float %b) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ftz_param_0];
-; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; 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;
diff --git a/llvm/test/CodeGen/NVPTX/sqrt-approx.ll b/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
index a60ab717f3755..a28d264cd8ec0 100644
--- a/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
+++ b/llvm/test/CodeGen/NVPTX/sqrt-approx.ll
@@ -83,7 +83,7 @@ define float @test_sqrt32(float %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_param_0];
-; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
@@ -115,7 +115,7 @@ define float @test_sqrt_ftz(float %a) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_ftz_param_0];
-; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
@@ -240,7 +240,7 @@ define float @test_sqrt32_refined(float %a) #0 #2 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_param_0];
-; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;
+; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
@@ -352,7 +352,7 @@ define float @test_sqrt32_refined_ftz(float %a) #0 #1 #2 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ftz_param_0];
-; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;
+; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%ret = tail call float @llvm.sqrt.f32(float %a)
More information about the llvm-commits
mailing list