test/CodeGen/NVPTX/fast-math.ll

; 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)

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
}

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
}

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
}

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
}

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
}

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
}

; 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.

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
}

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
}

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
}

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
}

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
}

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
}

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
}

declare float @llvm.sin.f32(float)
declare float @llvm.cos.f32(float)

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
}

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
}

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
}

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
}

define float @repeated_div_recip_allowed(i1 %pred, float %a, float %b, float %divisor) {
; 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
  %w = select i1 %pred, float %z, float %y
  ret float %w
}

define float @repeated_div_recip_allowed_sel(i1 %pred, float %a, float %b, float %divisor) {
; 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
}

define float @repeated_div_recip_allowed_ftz(i1 %pred, float %a, float %b, float %divisor) #1 {
; 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
  %w = select i1 %pred, float %z, float %y
  ret float %w
}

define float @repeated_div_recip_allowed_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #1 {
; 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
}

define float @repeated_div_fast(i1 %pred, float %a, float %b, float %divisor) #0 {
; 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
  %w = select i1 %pred, float %z, float %y
  ret float %w
}

define float @repeated_div_fast_sel(i1 %pred, float %a, float %b, float %divisor) #0 {
; 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
}

define float @repeated_div_fast_ftz(i1 %pred, float %a, float %b, float %divisor) #0 #1 {
; 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
  %w = select i1 %pred, float %z, float %y
  ret float %w
}

define float @repeated_div_fast_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #0 #1 {
; 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
}

define float @frem(float %a, float %b) #0 {
; 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
}

define float @frem_ftz(float %a, float %b) #0 #1 {
; 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
}

define double @frem_f64(double %a, double %b) #0 {
; 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
}

attributes #0 = { "unsafe-fp-math" = "true" }
attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }