441 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 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \3; RUN: | FileCheck %s4; RUN: %if ptxas %{ \5; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \6; RUN: | %ptxas-verify \7; RUN: %}8 9target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"10 11declare float @llvm.sqrt.f32(float)12declare double @llvm.sqrt.f64(double)13 14; -- reciprocal sqrt --15 16define float @test_rsqrt32(float %a) {17; CHECK-LABEL: test_rsqrt32(18; CHECK: {19; CHECK-NEXT: .reg .b32 %r<3>;20; CHECK-EMPTY:21; CHECK-NEXT: // %bb.0:22; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_param_0];23; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;24; CHECK-NEXT: st.param.b32 [func_retval0], %r2;25; CHECK-NEXT: ret;26 %val = tail call float @llvm.sqrt.f32(float %a)27 %ret = fdiv float 1.0, %val28 ret float %ret29}30 31define float @test_rsqrt_ftz(float %a) #1 {32; CHECK-LABEL: test_rsqrt_ftz(33; CHECK: {34; CHECK-NEXT: .reg .b32 %r<3>;35; CHECK-EMPTY:36; CHECK-NEXT: // %bb.0:37; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt_ftz_param_0];38; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;39; CHECK-NEXT: st.param.b32 [func_retval0], %r2;40; CHECK-NEXT: ret;41 %val = tail call float @llvm.sqrt.f32(float %a)42 %ret = fdiv float 1.0, %val43 ret float %ret44}45 46define double @test_rsqrt64(double %a) {47; CHECK-LABEL: test_rsqrt64(48; CHECK: {49; CHECK-NEXT: .reg .b64 %rd<3>;50; CHECK-EMPTY:51; CHECK-NEXT: // %bb.0:52; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_param_0];53; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;54; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;55; CHECK-NEXT: ret;56 %val = tail call double @llvm.sqrt.f64(double %a)57 %ret = fdiv arcp double 1.0, %val58 ret double %ret59}60 61; There's no rsqrt.approx.ftz.f64 instruction; we just use the non-ftz version.62define double @test_rsqrt64_ftz(double %a) #1 {63; CHECK-LABEL: test_rsqrt64_ftz(64; CHECK: {65; CHECK-NEXT: .reg .b64 %rd<3>;66; CHECK-EMPTY:67; CHECK-NEXT: // %bb.0:68; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_ftz_param_0];69; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;70; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;71; CHECK-NEXT: ret;72 %val = tail call double @llvm.sqrt.f64(double %a)73 %ret = fdiv arcp double 1.0, %val74 ret double %ret75}76 77; -- sqrt --78 79define float @test_sqrt32(float %a) {80; CHECK-LABEL: test_sqrt32(81; CHECK: {82; CHECK-NEXT: .reg .b32 %r<3>;83; CHECK-EMPTY:84; CHECK-NEXT: // %bb.0:85; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_param_0];86; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;87; CHECK-NEXT: st.param.b32 [func_retval0], %r2;88; CHECK-NEXT: ret;89 %ret = tail call float @llvm.sqrt.f32(float %a)90 ret float %ret91}92 93define float @test_sqrt32_ninf(float %a) {94; CHECK-LABEL: test_sqrt32_ninf(95; CHECK: {96; CHECK-NEXT: .reg .pred %p<2>;97; CHECK-NEXT: .reg .b32 %r<5>;98; CHECK-EMPTY:99; CHECK-NEXT: // %bb.0:100; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_ninf_param_0];101; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;102; CHECK-NEXT: abs.f32 %r3, %r1;103; CHECK-NEXT: setp.lt.f32 %p1, %r3, 0f00800000;104; CHECK-NEXT: selp.f32 %r4, 0f00000000, %r2, %p1;105; CHECK-NEXT: st.param.b32 [func_retval0], %r4;106; CHECK-NEXT: ret;107 %ret = tail call ninf afn float @llvm.sqrt.f32(float %a)108 ret float %ret109}110 111define float @test_sqrt_ftz(float %a) #1 {112; CHECK-LABEL: test_sqrt_ftz(113; CHECK: {114; CHECK-NEXT: .reg .b32 %r<3>;115; CHECK-EMPTY:116; CHECK-NEXT: // %bb.0:117; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_ftz_param_0];118; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;119; CHECK-NEXT: st.param.b32 [func_retval0], %r2;120; CHECK-NEXT: ret;121 %ret = tail call float @llvm.sqrt.f32(float %a)122 ret float %ret123}124 125define float @test_sqrt_ftz_ninf(float %a) #1 {126; CHECK-LABEL: test_sqrt_ftz_ninf(127; CHECK: {128; CHECK-NEXT: .reg .pred %p<2>;129; CHECK-NEXT: .reg .b32 %r<4>;130; CHECK-EMPTY:131; CHECK-NEXT: // %bb.0:132; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_ftz_ninf_param_0];133; CHECK-NEXT: setp.eq.ftz.f32 %p1, %r1, 0f00000000;134; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;135; CHECK-NEXT: selp.f32 %r3, 0f00000000, %r2, %p1;136; CHECK-NEXT: st.param.b32 [func_retval0], %r3;137; CHECK-NEXT: ret;138 %ret = tail call ninf afn float @llvm.sqrt.f32(float %a)139 ret float %ret140}141 142define double @test_sqrt64(double %a) {143; CHECK-LABEL: test_sqrt64(144; CHECK: {145; CHECK-NEXT: .reg .b64 %rd<3>;146; CHECK-EMPTY:147; CHECK-NEXT: // %bb.0:148; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_param_0];149; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;150; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;151; CHECK-NEXT: ret;152 %ret = tail call double @llvm.sqrt.f64(double %a)153 ret double %ret154}155 156; There's no sqrt.approx.f64 instruction; we emit157; reciprocal(rsqrt.approx.f64(x)). There's no non-ftz approximate reciprocal,158; so we just use the ftz version.159define double @test_sqrt64_ninf(double %a) {160; CHECK-LABEL: test_sqrt64_ninf(161; CHECK: {162; CHECK-NEXT: .reg .pred %p<2>;163; CHECK-NEXT: .reg .b64 %rd<6>;164; CHECK-EMPTY:165; CHECK-NEXT: // %bb.0:166; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_ninf_param_0];167; CHECK-NEXT: abs.f64 %rd2, %rd1;168; CHECK-NEXT: setp.lt.f64 %p1, %rd2, 0d0010000000000000;169; CHECK-NEXT: rsqrt.approx.f64 %rd3, %rd1;170; CHECK-NEXT: rcp.approx.ftz.f64 %rd4, %rd3;171; CHECK-NEXT: selp.f64 %rd5, 0d0000000000000000, %rd4, %p1;172; CHECK-NEXT: st.param.b64 [func_retval0], %rd5;173; CHECK-NEXT: ret;174 %ret = tail call ninf afn double @llvm.sqrt.f64(double %a)175 ret double %ret176}177 178define double @test_sqrt64_ftz(double %a) #1 {179; CHECK-LABEL: test_sqrt64_ftz(180; CHECK: {181; CHECK-NEXT: .reg .b64 %rd<3>;182; CHECK-EMPTY:183; CHECK-NEXT: // %bb.0:184; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_ftz_param_0];185; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;186; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;187; CHECK-NEXT: ret;188 %ret = tail call double @llvm.sqrt.f64(double %a)189 ret double %ret190}191 192; There's no sqrt.approx.ftz.f64 instruction; we just use the non-ftz version.193define double @test_sqrt64_ftz_ninf(double %a) #1 {194; CHECK-LABEL: test_sqrt64_ftz_ninf(195; CHECK: {196; CHECK-NEXT: .reg .pred %p<2>;197; CHECK-NEXT: .reg .b64 %rd<6>;198; CHECK-EMPTY:199; CHECK-NEXT: // %bb.0:200; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_ftz_ninf_param_0];201; CHECK-NEXT: abs.f64 %rd2, %rd1;202; CHECK-NEXT: setp.lt.f64 %p1, %rd2, 0d0010000000000000;203; CHECK-NEXT: rsqrt.approx.f64 %rd3, %rd1;204; CHECK-NEXT: rcp.approx.ftz.f64 %rd4, %rd3;205; CHECK-NEXT: selp.f64 %rd5, 0d0000000000000000, %rd4, %p1;206; CHECK-NEXT: st.param.b64 [func_retval0], %rd5;207; CHECK-NEXT: ret;208 %ret = tail call ninf afn double @llvm.sqrt.f64(double %a)209 ret double %ret210}211 212; -- refined sqrt and rsqrt --213;214; The sqrt and rsqrt refinement algorithms both emit an rsqrt.approx, followed215; by some math.216 217define float @test_rsqrt32_refined(float %a) #2 {218; CHECK-LABEL: test_rsqrt32_refined(219; CHECK: {220; CHECK-NEXT: .reg .b32 %r<7>;221; CHECK-EMPTY:222; CHECK-NEXT: // %bb.0:223; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_refined_param_0];224; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;225; CHECK-NEXT: mul.f32 %r3, %r1, %r2;226; CHECK-NEXT: fma.rn.f32 %r4, %r3, %r2, 0fC0400000;227; CHECK-NEXT: mul.f32 %r5, %r2, 0fBF000000;228; CHECK-NEXT: mul.f32 %r6, %r5, %r4;229; CHECK-NEXT: st.param.b32 [func_retval0], %r6;230; CHECK-NEXT: ret;231 %val = tail call float @llvm.sqrt.f32(float %a)232 %ret = fdiv arcp contract float 1.0, %val233 ret float %ret234}235 236define float @test_sqrt32_refined(float %a) #2 {237; CHECK-LABEL: test_sqrt32_refined(238; CHECK: {239; CHECK-NEXT: .reg .b32 %r<3>;240; CHECK-EMPTY:241; CHECK-NEXT: // %bb.0:242; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_param_0];243; CHECK-NEXT: sqrt.approx.f32 %r2, %r1;244; CHECK-NEXT: st.param.b32 [func_retval0], %r2;245; CHECK-NEXT: ret;246 %ret = tail call float @llvm.sqrt.f32(float %a)247 ret float %ret248}249 250define float @test_sqrt32_refined_ninf(float %a) #2 {251; CHECK-LABEL: test_sqrt32_refined_ninf(252; CHECK: {253; CHECK-NEXT: .reg .pred %p<2>;254; CHECK-NEXT: .reg .b32 %r<9>;255; CHECK-EMPTY:256; CHECK-NEXT: // %bb.0:257; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ninf_param_0];258; CHECK-NEXT: rsqrt.approx.f32 %r2, %r1;259; CHECK-NEXT: mul.f32 %r3, %r1, %r2;260; CHECK-NEXT: fma.rn.f32 %r4, %r3, %r2, 0fC0400000;261; CHECK-NEXT: mul.f32 %r5, %r3, 0fBF000000;262; CHECK-NEXT: mul.f32 %r6, %r5, %r4;263; CHECK-NEXT: abs.f32 %r7, %r1;264; CHECK-NEXT: setp.lt.f32 %p1, %r7, 0f00800000;265; CHECK-NEXT: selp.f32 %r8, 0f00000000, %r6, %p1;266; CHECK-NEXT: st.param.b32 [func_retval0], %r8;267; CHECK-NEXT: ret;268 %ret = tail call ninf afn contract float @llvm.sqrt.f32(float %a)269 ret float %ret270}271 272define double @test_rsqrt64_refined(double %a) #2 {273; CHECK-LABEL: test_rsqrt64_refined(274; CHECK: {275; CHECK-NEXT: .reg .b64 %rd<7>;276; CHECK-EMPTY:277; CHECK-NEXT: // %bb.0:278; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_refined_param_0];279; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;280; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;281; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;282; CHECK-NEXT: mul.f64 %rd5, %rd2, 0dBFE0000000000000;283; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;284; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;285; CHECK-NEXT: ret;286 %val = tail call double @llvm.sqrt.f64(double %a)287 %ret = fdiv arcp contract double 1.0, %val288 ret double %ret289}290 291define double @test_sqrt64_refined(double %a) #2 {292; CHECK-LABEL: test_sqrt64_refined(293; CHECK: {294; CHECK-NEXT: .reg .b64 %rd<3>;295; CHECK-EMPTY:296; CHECK-NEXT: // %bb.0:297; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_param_0];298; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;299; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;300; CHECK-NEXT: ret;301 %ret = tail call double @llvm.sqrt.f64(double %a)302 ret double %ret303}304 305define double @test_sqrt64_refined_ninf(double %a) #2 {306; CHECK-LABEL: test_sqrt64_refined_ninf(307; CHECK: {308; CHECK-NEXT: .reg .pred %p<2>;309; CHECK-NEXT: .reg .b64 %rd<9>;310; CHECK-EMPTY:311; CHECK-NEXT: // %bb.0:312; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ninf_param_0];313; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;314; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;315; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;316; CHECK-NEXT: mul.f64 %rd5, %rd3, 0dBFE0000000000000;317; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;318; CHECK-NEXT: abs.f64 %rd7, %rd1;319; CHECK-NEXT: setp.lt.f64 %p1, %rd7, 0d0010000000000000;320; CHECK-NEXT: selp.f64 %rd8, 0d0000000000000000, %rd6, %p1;321; CHECK-NEXT: st.param.b64 [func_retval0], %rd8;322; CHECK-NEXT: ret;323 %ret = tail call ninf afn contract double @llvm.sqrt.f64(double %a)324 ret double %ret325}326 327; -- refined sqrt and rsqrt with ftz enabled --328 329define float @test_rsqrt32_refined_ftz(float %a) #1 #2 {330; CHECK-LABEL: test_rsqrt32_refined_ftz(331; CHECK: {332; CHECK-NEXT: .reg .b32 %r<7>;333; CHECK-EMPTY:334; CHECK-NEXT: // %bb.0:335; CHECK-NEXT: ld.param.b32 %r1, [test_rsqrt32_refined_ftz_param_0];336; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;337; CHECK-NEXT: mul.ftz.f32 %r3, %r1, %r2;338; CHECK-NEXT: fma.rn.ftz.f32 %r4, %r3, %r2, 0fC0400000;339; CHECK-NEXT: mul.ftz.f32 %r5, %r2, 0fBF000000;340; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r4;341; CHECK-NEXT: st.param.b32 [func_retval0], %r6;342; CHECK-NEXT: ret;343 %val = tail call float @llvm.sqrt.f32(float %a)344 %ret = fdiv arcp contract float 1.0, %val345 ret float %ret346}347 348define float @test_sqrt32_refined_ftz(float %a) #1 #2 {349; CHECK-LABEL: test_sqrt32_refined_ftz(350; CHECK: {351; CHECK-NEXT: .reg .b32 %r<3>;352; CHECK-EMPTY:353; CHECK-NEXT: // %bb.0:354; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ftz_param_0];355; CHECK-NEXT: sqrt.approx.ftz.f32 %r2, %r1;356; CHECK-NEXT: st.param.b32 [func_retval0], %r2;357; CHECK-NEXT: ret;358 %ret = tail call float @llvm.sqrt.f32(float %a)359 ret float %ret360}361 362define float @test_sqrt32_refined_ftz_ninf(float %a) #1 #2 {363; CHECK-LABEL: test_sqrt32_refined_ftz_ninf(364; CHECK: {365; CHECK-NEXT: .reg .pred %p<2>;366; CHECK-NEXT: .reg .b32 %r<8>;367; CHECK-EMPTY:368; CHECK-NEXT: // %bb.0:369; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt32_refined_ftz_ninf_param_0];370; CHECK-NEXT: rsqrt.approx.ftz.f32 %r2, %r1;371; CHECK-NEXT: mul.ftz.f32 %r3, %r1, %r2;372; CHECK-NEXT: fma.rn.ftz.f32 %r4, %r3, %r2, 0fC0400000;373; CHECK-NEXT: mul.ftz.f32 %r5, %r3, 0fBF000000;374; CHECK-NEXT: mul.ftz.f32 %r6, %r5, %r4;375; CHECK-NEXT: setp.eq.ftz.f32 %p1, %r1, 0f00000000;376; CHECK-NEXT: selp.f32 %r7, 0f00000000, %r6, %p1;377; CHECK-NEXT: st.param.b32 [func_retval0], %r7;378; CHECK-NEXT: ret;379 %ret = tail call ninf afn contract float @llvm.sqrt.f32(float %a)380 ret float %ret381}382 383; There's no rsqrt.approx.ftz.f64, so we just use the non-ftz version.384define double @test_rsqrt64_refined_ftz(double %a) #1 #2 {385; CHECK-LABEL: test_rsqrt64_refined_ftz(386; CHECK: {387; CHECK-NEXT: .reg .b64 %rd<7>;388; CHECK-EMPTY:389; CHECK-NEXT: // %bb.0:390; CHECK-NEXT: ld.param.b64 %rd1, [test_rsqrt64_refined_ftz_param_0];391; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;392; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;393; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;394; CHECK-NEXT: mul.f64 %rd5, %rd2, 0dBFE0000000000000;395; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;396; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;397; CHECK-NEXT: ret;398 %val = tail call double @llvm.sqrt.f64(double %a)399 %ret = fdiv arcp contract double 1.0, %val400 ret double %ret401}402 403define double @test_sqrt64_refined_ftz(double %a) #1 #2 {404; CHECK-LABEL: test_sqrt64_refined_ftz(405; CHECK: {406; CHECK-NEXT: .reg .b64 %rd<3>;407; CHECK-EMPTY:408; CHECK-NEXT: // %bb.0:409; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ftz_param_0];410; CHECK-NEXT: sqrt.rn.f64 %rd2, %rd1;411; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;412; CHECK-NEXT: ret;413 %ret = tail call double @llvm.sqrt.f64(double %a)414 ret double %ret415}416 417define double @test_sqrt64_refined_ftz_ninf(double %a) #1 #2 {418; CHECK-LABEL: test_sqrt64_refined_ftz_ninf(419; CHECK: {420; CHECK-NEXT: .reg .pred %p<2>;421; CHECK-NEXT: .reg .b64 %rd<9>;422; CHECK-EMPTY:423; CHECK-NEXT: // %bb.0:424; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt64_refined_ftz_ninf_param_0];425; CHECK-NEXT: rsqrt.approx.f64 %rd2, %rd1;426; CHECK-NEXT: mul.f64 %rd3, %rd1, %rd2;427; CHECK-NEXT: fma.rn.f64 %rd4, %rd3, %rd2, 0dC008000000000000;428; CHECK-NEXT: mul.f64 %rd5, %rd3, 0dBFE0000000000000;429; CHECK-NEXT: mul.f64 %rd6, %rd5, %rd4;430; CHECK-NEXT: abs.f64 %rd7, %rd1;431; CHECK-NEXT: setp.lt.f64 %p1, %rd7, 0d0010000000000000;432; CHECK-NEXT: selp.f64 %rd8, 0d0000000000000000, %rd6, %p1;433; CHECK-NEXT: st.param.b64 [func_retval0], %rd8;434; CHECK-NEXT: ret;435 %ret = tail call ninf afn contract double @llvm.sqrt.f64(double %a)436 ret double %ret437}438 439attributes #1 = { "denormal-fp-math-f32" = "preserve-sign,preserve-sign" }440attributes #2 = { "reciprocal-estimates" = "rsqrtf:1,rsqrtd:1,sqrtf:1,sqrtd:1" }441