553 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5declare float @llvm.sqrt.f32(float)6declare double @llvm.sqrt.f64(double)7 8define float @sqrt_div(float %a, float %b) {9; CHECK-LABEL: sqrt_div(10; CHECK: {11; CHECK-NEXT: .reg .b32 %r<5>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_param_0];15; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;16; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_param_1];17; CHECK-NEXT: div.rn.f32 %r4, %r2, %r3;18; CHECK-NEXT: st.param.b32 [func_retval0], %r4;19; CHECK-NEXT: ret;20 %t1 = tail call float @llvm.sqrt.f32(float %a)21 %t2 = fdiv float %t1, %b22 ret float %t223}24 25define float @sqrt_div_fast(float %a, float %b) {26; CHECK-LABEL: sqrt_div_fast(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<5>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_param_0];32; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;33; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_param_1];34; CHECK-NEXT: div.approx.f32 %r4, %r2, %r3;35; CHECK-NEXT: st.param.b32 [func_retval0], %r4;36; CHECK-NEXT: ret;37 %t1 = tail call afn float @llvm.sqrt.f32(float %a)38 %t2 = fdiv afn float %t1, %b39 ret float %t240}41 42define float @sqrt_div_fast_ninf(float %a, float %b) {43; CHECK-LABEL: sqrt_div_fast_ninf(44; CHECK: {45; CHECK-NEXT: .reg .b32 %r<5>;46; CHECK-EMPTY:47; CHECK-NEXT: // %bb.0:48; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ninf_param_0];49; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;50; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_ninf_param_1];51; CHECK-NEXT: div.approx.f32 %r4, %r2, %r3;52; CHECK-NEXT: st.param.b32 [func_retval0], %r4;53; CHECK-NEXT: ret;54 %t1 = tail call ninf afn float @llvm.sqrt.f32(float %a)55 %t2 = fdiv afn float %t1, %b56 ret float %t257}58 59define float @sqrt_div_ftz(float %a, float %b) #1 {60; CHECK-LABEL: sqrt_div_ftz(61; CHECK: {62; CHECK-NEXT: .reg .b32 %r<5>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0:65; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_ftz_param_0];66; CHECK-NEXT: sqrt.rn.ftz.f32 %r2, %r1;67; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_ftz_param_1];68; CHECK-NEXT: div.rn.ftz.f32 %r4, %r2, %r3;69; CHECK-NEXT: st.param.b32 [func_retval0], %r4;70; CHECK-NEXT: ret;71 %t1 = tail call float @llvm.sqrt.f32(float %a)72 %t2 = fdiv float %t1, %b73 ret float %t274}75 76define float @sqrt_div_fast_ftz(float %a, float %b) #1 {77; CHECK-LABEL: sqrt_div_fast_ftz(78; CHECK: {79; CHECK-NEXT: .reg .b32 %r<5>;80; CHECK-EMPTY:81; CHECK-NEXT: // %bb.0:82; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ftz_param_0];83; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;84; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_ftz_param_1];85; CHECK-NEXT: div.approx.ftz.f32 %r4, %r2, %r3;86; CHECK-NEXT: st.param.b32 [func_retval0], %r4;87; CHECK-NEXT: ret;88 %t1 = tail call afn float @llvm.sqrt.f32(float %a)89 %t2 = fdiv afn float %t1, %b90 ret float %t291}92 93define float @sqrt_div_fast_ftz_ninf(float %a, float %b) #1 {94; CHECK-LABEL: sqrt_div_fast_ftz_ninf(95; CHECK: {96; CHECK-NEXT: .reg .b32 %r<5>;97; CHECK-EMPTY:98; CHECK-NEXT: // %bb.0:99; CHECK-NEXT: ld.param.b32 %r1, [sqrt_div_fast_ftz_ninf_param_0];100; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;101; CHECK-NEXT: ld.param.b32 %r3, [sqrt_div_fast_ftz_ninf_param_1];102; CHECK-NEXT: div.approx.ftz.f32 %r4, %r2, %r3;103; CHECK-NEXT: st.param.b32 [func_retval0], %r4;104; CHECK-NEXT: ret;105 %t1 = tail call ninf afn float @llvm.sqrt.f32(float %a)106 %t2 = fdiv afn float %t1, %b107 ret float %t2108}109 110; There are no fast-math or ftz versions of sqrt and div for f64. We use111; reciprocal(rsqrt(x)) for sqrt(x), and emit a vanilla divide.112 113define double @sqrt_div_fast_ftz_f64(double %a, double %b) #1 {114; CHECK-LABEL: sqrt_div_fast_ftz_f64(115; CHECK: {116; CHECK-NEXT: .reg .b64 %rd<5>;117; CHECK-EMPTY:118; CHECK-NEXT: // %bb.0:119; CHECK-NEXT: ld.param.b64 %rd1, [sqrt_div_fast_ftz_f64_param_0];120; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;121; CHECK-NEXT: ld.param.b64 %rd3, [sqrt_div_fast_ftz_f64_param_1];122; CHECK-NEXT: div.rn.f64 %rd4, %rd2, %rd3;123; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;124; CHECK-NEXT: ret;125 %t1 = tail call double @llvm.sqrt.f64(double %a)126 %t2 = fdiv double %t1, %b127 ret double %t2128}129 130define double @sqrt_div_fast_ftz_f64_ninf(double %a, double %b) #1 {131; CHECK-LABEL: sqrt_div_fast_ftz_f64_ninf(132; CHECK: {133; CHECK-NEXT: .reg .b64 %rd<5>;134; CHECK-EMPTY:135; CHECK-NEXT: // %bb.0:136; CHECK-NEXT: ld.param.b64 %rd1, [sqrt_div_fast_ftz_f64_ninf_param_0];137; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;138; CHECK-NEXT: ld.param.b64 %rd3, [sqrt_div_fast_ftz_f64_ninf_param_1];139; CHECK-NEXT: div.rn.f64 %rd4, %rd2, %rd3;140; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;141; CHECK-NEXT: ret;142 %t1 = tail call ninf afn double @llvm.sqrt.f64(double %a)143 %t2 = fdiv double %t1, %b144 ret double %t2145}146 147define float @rsqrt(float %a) {148; CHECK-LABEL: rsqrt(149; CHECK: {150; CHECK-NEXT: .reg .b32 %r<4>;151; CHECK-EMPTY:152; CHECK-NEXT: // %bb.0:153; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_param_0];154; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;155; CHECK-NEXT: rcp.rn.f32 %r3, %r2;156; CHECK-NEXT: st.param.b32 [func_retval0], %r3;157; CHECK-NEXT: ret;158 %b = tail call float @llvm.sqrt.f32(float %a)159 %ret = fdiv float 1.0, %b160 ret float %ret161}162 163define float @rsqrt_fast(float %a) {164; CHECK-LABEL: rsqrt_fast(165; CHECK: {166; CHECK-NEXT: .reg .b32 %r<3>;167; CHECK-EMPTY:168; CHECK-NEXT: // %bb.0:169; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_fast_param_0];170; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;171; CHECK-NEXT: st.param.b32 [func_retval0], %r2;172; CHECK-NEXT: ret;173 %b = tail call afn float @llvm.sqrt.f32(float %a)174 %ret = fdiv afn float 1.0, %b175 ret float %ret176}177 178define float @rsqrt_fast_ftz(float %a) #1 {179; CHECK-LABEL: rsqrt_fast_ftz(180; CHECK: {181; CHECK-NEXT: .reg .b32 %r<3>;182; CHECK-EMPTY:183; CHECK-NEXT: // %bb.0:184; CHECK-NEXT: ld.param.b32 %r1, [rsqrt_fast_ftz_param_0];185; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;186; CHECK-NEXT: st.param.b32 [func_retval0], %r2;187; CHECK-NEXT: ret;188 %b = tail call afn float @llvm.sqrt.f32(float %a)189 %ret = fdiv afn float 1.0, %b190 ret float %ret191}192 193define float @fadd(float %a, float %b) {194; CHECK-LABEL: fadd(195; CHECK: {196; CHECK-NEXT: .reg .b32 %r<4>;197; CHECK-EMPTY:198; CHECK-NEXT: // %bb.0:199; CHECK-NEXT: ld.param.b32 %r1, [fadd_param_0];200; CHECK-NEXT: ld.param.b32 %r2, [fadd_param_1];201; CHECK-NEXT: add.rn.f32 %r3, %r1, %r2;202; CHECK-NEXT: st.param.b32 [func_retval0], %r3;203; CHECK-NEXT: ret;204 %t1 = fadd float %a, %b205 ret float %t1206}207 208define float @fadd_ftz(float %a, float %b) #1 {209; CHECK-LABEL: fadd_ftz(210; CHECK: {211; CHECK-NEXT: .reg .b32 %r<4>;212; CHECK-EMPTY:213; CHECK-NEXT: // %bb.0:214; CHECK-NEXT: ld.param.b32 %r1, [fadd_ftz_param_0];215; CHECK-NEXT: ld.param.b32 %r2, [fadd_ftz_param_1];216; CHECK-NEXT: add.rn.ftz.f32 %r3, %r1, %r2;217; CHECK-NEXT: st.param.b32 [func_retval0], %r3;218; CHECK-NEXT: ret;219 %t1 = fadd float %a, %b220 ret float %t1221}222 223declare float @llvm.sin.f32(float)224declare float @llvm.cos.f32(float)225 226define float @fsin_approx_afn(float %a) {227; CHECK-LABEL: fsin_approx_afn(228; CHECK: {229; CHECK-NEXT: .reg .b32 %r<3>;230; CHECK-EMPTY:231; CHECK-NEXT: // %bb.0:232; CHECK-NEXT: ld.param.b32 %r1, [fsin_approx_afn_param_0];233; CHECK-NEXT: sin.approx.f32 %r2, %r1;234; CHECK-NEXT: st.param.b32 [func_retval0], %r2;235; CHECK-NEXT: ret;236 %r = tail call afn float @llvm.sin.f32(float %a)237 ret float %r238}239 240define float @fcos_approx_afn(float %a) {241; CHECK-LABEL: fcos_approx_afn(242; CHECK: {243; CHECK-NEXT: .reg .b32 %r<3>;244; CHECK-EMPTY:245; CHECK-NEXT: // %bb.0:246; CHECK-NEXT: ld.param.b32 %r1, [fcos_approx_afn_param_0];247; CHECK-NEXT: cos.approx.f32 %r2, %r1;248; CHECK-NEXT: st.param.b32 [func_retval0], %r2;249; CHECK-NEXT: ret;250 %r = tail call afn float @llvm.cos.f32(float %a)251 ret float %r252}253 254define float @fsin_approx_ftz(float %a) #1 {255; CHECK-LABEL: fsin_approx_ftz(256; CHECK: {257; CHECK-NEXT: .reg .b32 %r<3>;258; CHECK-EMPTY:259; CHECK-NEXT: // %bb.0:260; CHECK-NEXT: ld.param.b32 %r1, [fsin_approx_ftz_param_0];261; CHECK-NEXT: sin.approx.ftz.f32 %r2, %r1;262; CHECK-NEXT: st.param.b32 [func_retval0], %r2;263; CHECK-NEXT: ret;264 %r = tail call afn float @llvm.sin.f32(float %a)265 ret float %r266}267 268define float @fcos_approx_ftz(float %a) #1 {269; CHECK-LABEL: fcos_approx_ftz(270; CHECK: {271; CHECK-NEXT: .reg .b32 %r<3>;272; CHECK-EMPTY:273; CHECK-NEXT: // %bb.0:274; CHECK-NEXT: ld.param.b32 %r1, [fcos_approx_ftz_param_0];275; CHECK-NEXT: cos.approx.ftz.f32 %r2, %r1;276; CHECK-NEXT: st.param.b32 [func_retval0], %r2;277; CHECK-NEXT: ret;278 %r = tail call afn float @llvm.cos.f32(float %a)279 ret float %r280}281 282define float @repeated_div_recip_allowed(i1 %pred, float %a, float %b, float %divisor) {283; CHECK-LABEL: repeated_div_recip_allowed(284; CHECK: {285; CHECK-NEXT: .reg .pred %p<2>;286; CHECK-NEXT: .reg .b16 %rs<3>;287; CHECK-NEXT: .reg .b32 %r<9>;288; CHECK-EMPTY:289; CHECK-NEXT: // %bb.0:290; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_param_0];291; CHECK-NEXT: and.b16 %rs2, %rs1, 1;292; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;293; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_param_1];294; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_param_3];295; CHECK-NEXT: rcp.rn.f32 %r3, %r2;296; CHECK-NEXT: mul.rn.f32 %r4, %r1, %r3;297; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_recip_allowed_param_2];298; CHECK-NEXT: mul.rn.f32 %r6, %r5, %r3;299; CHECK-NEXT: mul.rn.f32 %r7, %r4, %r6;300; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p1;301; CHECK-NEXT: st.param.b32 [func_retval0], %r8;302; CHECK-NEXT: ret;303 %x = fdiv arcp float %a, %divisor304 %y = fdiv arcp float %b, %divisor305 %z = fmul float %x, %y306 %w = select i1 %pred, float %z, float %y307 ret float %w308}309 310define float @repeated_div_recip_allowed_sel(i1 %pred, float %a, float %b, float %divisor) {311; CHECK-LABEL: repeated_div_recip_allowed_sel(312; CHECK: {313; CHECK-NEXT: .reg .pred %p<2>;314; CHECK-NEXT: .reg .b16 %rs<3>;315; CHECK-NEXT: .reg .b32 %r<4>;316; CHECK-NEXT: .reg .b64 %rd<4>;317; CHECK-EMPTY:318; CHECK-NEXT: // %bb.0:319; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_sel_param_0];320; CHECK-NEXT: and.b16 %rs2, %rs1, 1;321; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;322; CHECK-NEXT: mov.b64 %rd1, repeated_div_recip_allowed_sel_param_2;323; CHECK-NEXT: mov.b64 %rd2, repeated_div_recip_allowed_sel_param_1;324; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;325; CHECK-NEXT: ld.param.b32 %r1, [%rd3];326; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_sel_param_3];327; CHECK-NEXT: div.rn.f32 %r3, %r1, %r2;328; CHECK-NEXT: st.param.b32 [func_retval0], %r3;329; CHECK-NEXT: ret;330 %x = fdiv arcp float %a, %divisor331 %y = fdiv arcp float %b, %divisor332 %w = select i1 %pred, float %x, float %y333 ret float %w334}335 336define float @repeated_div_recip_allowed_ftz(i1 %pred, float %a, float %b, float %divisor) #1 {337; CHECK-LABEL: repeated_div_recip_allowed_ftz(338; CHECK: {339; CHECK-NEXT: .reg .pred %p<2>;340; CHECK-NEXT: .reg .b16 %rs<3>;341; CHECK-NEXT: .reg .b32 %r<9>;342; CHECK-EMPTY:343; CHECK-NEXT: // %bb.0:344; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_ftz_param_0];345; CHECK-NEXT: and.b16 %rs2, %rs1, 1;346; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;347; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_ftz_param_1];348; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_param_3];349; CHECK-NEXT: rcp.rn.ftz.f32 %r3, %r2;350; CHECK-NEXT: mul.rn.ftz.f32 %r4, %r1, %r3;351; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_recip_allowed_ftz_param_2];352; CHECK-NEXT: mul.rn.ftz.f32 %r6, %r5, %r3;353; CHECK-NEXT: mul.rn.ftz.f32 %r7, %r4, %r6;354; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p1;355; CHECK-NEXT: st.param.b32 [func_retval0], %r8;356; CHECK-NEXT: ret;357 %x = fdiv arcp float %a, %divisor358 %y = fdiv arcp float %b, %divisor359 %z = fmul float %x, %y360 %w = select i1 %pred, float %z, float %y361 ret float %w362}363 364define float @repeated_div_recip_allowed_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #1 {365; CHECK-LABEL: repeated_div_recip_allowed_ftz_sel(366; CHECK: {367; CHECK-NEXT: .reg .pred %p<2>;368; CHECK-NEXT: .reg .b16 %rs<3>;369; CHECK-NEXT: .reg .b32 %r<4>;370; CHECK-NEXT: .reg .b64 %rd<4>;371; CHECK-EMPTY:372; CHECK-NEXT: // %bb.0:373; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_ftz_sel_param_0];374; CHECK-NEXT: and.b16 %rs2, %rs1, 1;375; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;376; CHECK-NEXT: mov.b64 %rd1, repeated_div_recip_allowed_ftz_sel_param_2;377; CHECK-NEXT: mov.b64 %rd2, repeated_div_recip_allowed_ftz_sel_param_1;378; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;379; CHECK-NEXT: ld.param.b32 %r1, [%rd3];380; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_sel_param_3];381; CHECK-NEXT: div.rn.ftz.f32 %r3, %r1, %r2;382; CHECK-NEXT: st.param.b32 [func_retval0], %r3;383; CHECK-NEXT: ret;384 %x = fdiv arcp float %a, %divisor385 %y = fdiv arcp float %b, %divisor386 %w = select i1 %pred, float %x, float %y387 ret float %w388}389 390define float @repeated_div_fast(i1 %pred, float %a, float %b, float %divisor) {391; CHECK-LABEL: repeated_div_fast(392; CHECK: {393; CHECK-NEXT: .reg .pred %p<2>;394; CHECK-NEXT: .reg .b16 %rs<3>;395; CHECK-NEXT: .reg .b32 %r<9>;396; CHECK-EMPTY:397; CHECK-NEXT: // %bb.0:398; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_param_0];399; CHECK-NEXT: and.b16 %rs2, %rs1, 1;400; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;401; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_fast_param_1];402; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_param_3];403; CHECK-NEXT: rcp.approx.f32 %r3, %r2;404; CHECK-NEXT: mul.f32 %r4, %r1, %r3;405; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_fast_param_2];406; CHECK-NEXT: mul.f32 %r6, %r5, %r3;407; CHECK-NEXT: mul.f32 %r7, %r4, %r6;408; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p1;409; CHECK-NEXT: st.param.b32 [func_retval0], %r8;410; CHECK-NEXT: ret;411 %x = fdiv afn arcp float %a, %divisor412 %y = fdiv afn arcp contract float %b, %divisor413 %z = fmul contract float %x, %y414 %w = select i1 %pred, float %z, float %y415 ret float %w416}417 418define float @repeated_div_fast_sel(i1 %pred, float %a, float %b, float %divisor) {419; CHECK-LABEL: repeated_div_fast_sel(420; CHECK: {421; CHECK-NEXT: .reg .pred %p<2>;422; CHECK-NEXT: .reg .b16 %rs<3>;423; CHECK-NEXT: .reg .b32 %r<4>;424; CHECK-NEXT: .reg .b64 %rd<4>;425; CHECK-EMPTY:426; CHECK-NEXT: // %bb.0:427; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_sel_param_0];428; CHECK-NEXT: and.b16 %rs2, %rs1, 1;429; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;430; CHECK-NEXT: mov.b64 %rd1, repeated_div_fast_sel_param_2;431; CHECK-NEXT: mov.b64 %rd2, repeated_div_fast_sel_param_1;432; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;433; CHECK-NEXT: ld.param.b32 %r1, [%rd3];434; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_sel_param_3];435; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;436; CHECK-NEXT: st.param.b32 [func_retval0], %r3;437; CHECK-NEXT: ret;438 %x = fdiv afn float %a, %divisor439 %y = fdiv afn float %b, %divisor440 %w = select i1 %pred, float %x, float %y441 ret float %w442}443 444define float @repeated_div_fast_ftz(i1 %pred, float %a, float %b, float %divisor) #1 {445; CHECK-LABEL: repeated_div_fast_ftz(446; CHECK: {447; CHECK-NEXT: .reg .pred %p<2>;448; CHECK-NEXT: .reg .b16 %rs<3>;449; CHECK-NEXT: .reg .b32 %r<9>;450; CHECK-EMPTY:451; CHECK-NEXT: // %bb.0:452; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_ftz_param_0];453; CHECK-NEXT: and.b16 %rs2, %rs1, 1;454; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;455; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_fast_ftz_param_1];456; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_param_3];457; CHECK-NEXT: rcp.approx.ftz.f32 %r3, %r2;458; CHECK-NEXT: mul.ftz.f32 %r4, %r1, %r3;459; CHECK-NEXT: ld.param.b32 %r5, [repeated_div_fast_ftz_param_2];460; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r3;461; CHECK-NEXT: mul.ftz.f32 %r7, %r4, %r6;462; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p1;463; CHECK-NEXT: st.param.b32 [func_retval0], %r8;464; CHECK-NEXT: ret;465 %x = fdiv afn arcp float %a, %divisor466 %y = fdiv afn arcp contract float %b, %divisor467 %z = fmul contract float %x, %y468 %w = select i1 %pred, float %z, float %y469 ret float %w470}471 472define float @repeated_div_fast_ftz_sel(i1 %pred, float %a, float %b, float %divisor) #1 {473; CHECK-LABEL: repeated_div_fast_ftz_sel(474; CHECK: {475; CHECK-NEXT: .reg .pred %p<2>;476; CHECK-NEXT: .reg .b16 %rs<3>;477; CHECK-NEXT: .reg .b32 %r<4>;478; CHECK-NEXT: .reg .b64 %rd<4>;479; CHECK-EMPTY:480; CHECK-NEXT: // %bb.0:481; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_ftz_sel_param_0];482; CHECK-NEXT: and.b16 %rs2, %rs1, 1;483; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;484; CHECK-NEXT: mov.b64 %rd1, repeated_div_fast_ftz_sel_param_2;485; CHECK-NEXT: mov.b64 %rd2, repeated_div_fast_ftz_sel_param_1;486; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;487; CHECK-NEXT: ld.param.b32 %r1, [%rd3];488; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_sel_param_3];489; CHECK-NEXT: div.approx.ftz.f32 %r3, %r1, %r2;490; CHECK-NEXT: st.param.b32 [func_retval0], %r3;491; CHECK-NEXT: ret;492 %x = fdiv afn float %a, %divisor493 %y = fdiv afn float %b, %divisor494 %w = select i1 %pred, float %x, float %y495 ret float %w496}497 498define float @frem(float %a, float %b) {499; CHECK-LABEL: frem(500; CHECK: {501; CHECK-NEXT: .reg .b32 %r<7>;502; CHECK-EMPTY:503; CHECK-NEXT: // %bb.0:504; CHECK-NEXT: ld.param.b32 %r1, [frem_param_0];505; CHECK-NEXT: ld.param.b32 %r2, [frem_param_1];506; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;507; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;508; CHECK-NEXT: neg.f32 %r5, %r4;509; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;510; CHECK-NEXT: st.param.b32 [func_retval0], %r6;511; CHECK-NEXT: ret;512 %rem = frem afn arcp contract ninf float %a, %b513 ret float %rem514}515 516define float @frem_ftz(float %a, float %b) #1 {517; CHECK-LABEL: frem_ftz(518; CHECK: {519; CHECK-NEXT: .reg .b32 %r<7>;520; CHECK-EMPTY:521; CHECK-NEXT: // %bb.0:522; CHECK-NEXT: ld.param.b32 %r1, [frem_ftz_param_0];523; CHECK-NEXT: ld.param.b32 %r2, [frem_ftz_param_1];524; CHECK-NEXT: div.approx.ftz.f32 %r3, %r1, %r2;525; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r4, %r3;526; CHECK-NEXT: neg.ftz.f32 %r5, %r4;527; CHECK-NEXT: fma.rn.ftz.f32 %r6, %r5, %r2, %r1;528; CHECK-NEXT: st.param.b32 [func_retval0], %r6;529; CHECK-NEXT: ret;530 %rem = frem afn contract ninf float %a, %b531 ret float %rem532}533 534define double @frem_f64(double %a, double %b) {535; CHECK-LABEL: frem_f64(536; CHECK: {537; CHECK-NEXT: .reg .b64 %rd<7>;538; CHECK-EMPTY:539; CHECK-NEXT: // %bb.0:540; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_param_0];541; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_param_1];542; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;543; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;544; CHECK-NEXT: neg.f64 %rd5, %rd4;545; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;546; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;547; CHECK-NEXT: ret;548 %rem = frem ninf double %a, %b549 ret double %rem550}551 552attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }553