167 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %}4 5define float @div_full(float %a, float %b) {6; CHECK-LABEL: div_full(7; CHECK: {8; CHECK-NEXT: .reg .b32 %r<9>;9; CHECK-EMPTY:10; CHECK-NEXT: // %bb.0:11; CHECK-NEXT: ld.param.b32 %r1, [div_full_param_0];12; CHECK-NEXT: ld.param.b32 %r2, [div_full_param_1];13; CHECK-NEXT: div.full.f32 %r3, %r1, %r2;14; CHECK-NEXT: mov.b32 %r4, 0f40400000;15; CHECK-NEXT: div.full.f32 %r5, %r3, %r4;16; CHECK-NEXT: div.full.ftz.f32 %r6, %r5, %r2;17; CHECK-NEXT: mov.b32 %r7, 0f40800000;18; CHECK-NEXT: div.full.ftz.f32 %r8, %r6, %r7;19; CHECK-NEXT: st.param.b32 [func_retval0], %r8;20; CHECK-NEXT: ret;21 %1 = call float @llvm.nvvm.div.full(float %a, float %b)22 %2 = call float @llvm.nvvm.div.full(float %1, float 3.0)23 %3 = call float @llvm.nvvm.div.full.ftz(float %2, float %b)24 %4 = call float @llvm.nvvm.div.full.ftz(float %3, float 4.0)25 ret float %426}27 28define float @div_fast_rr(float %a, float %b) {29; CHECK-LABEL: div_fast_rr(30; CHECK: {31; CHECK-NEXT: .reg .b32 %r<4>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b32 %r1, [div_fast_rr_param_0];35; CHECK-NEXT: ld.param.b32 %r2, [div_fast_rr_param_1];36; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;37; CHECK-NEXT: st.param.b32 [func_retval0], %r3;38; CHECK-NEXT: ret;39 %t1 = fdiv afn float %a, %b40 ret float %t141}42 43define float @div_fast_rr_ftz(float %a, float %b) #0 {44; CHECK-LABEL: div_fast_rr_ftz(45; CHECK: {46; CHECK-NEXT: .reg .b32 %r<4>;47; CHECK-EMPTY:48; CHECK-NEXT: // %bb.0:49; CHECK-NEXT: ld.param.b32 %r1, [div_fast_rr_ftz_param_0];50; CHECK-NEXT: ld.param.b32 %r2, [div_fast_rr_ftz_param_1];51; CHECK-NEXT: div.approx.ftz.f32 %r3, %r1, %r2;52; CHECK-NEXT: st.param.b32 [func_retval0], %r3;53; CHECK-NEXT: ret;54 %t1 = fdiv afn float %a, %b55 ret float %t156}57 58define float @div_fast_ri(float %a, float %b) {59; CHECK-LABEL: div_fast_ri(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<3>;62; CHECK-EMPTY:63; CHECK-NEXT: // %bb.0:64; CHECK-NEXT: ld.param.b32 %r1, [div_fast_ri_param_0];65; CHECK-NEXT: mul.rn.f32 %r2, %r1, 0f3F000000;66; CHECK-NEXT: st.param.b32 [func_retval0], %r2;67; CHECK-NEXT: ret;68 %t1 = fdiv afn float %a, 2.069 ret float %t170}71 72define float @div_fast_ri_ftz(float %a, float %b) #0 {73; CHECK-LABEL: div_fast_ri_ftz(74; CHECK: {75; CHECK-NEXT: .reg .b32 %r<3>;76; CHECK-EMPTY:77; CHECK-NEXT: // %bb.0:78; CHECK-NEXT: ld.param.b32 %r1, [div_fast_ri_ftz_param_0];79; CHECK-NEXT: mul.rn.ftz.f32 %r2, %r1, 0f3F000000;80; CHECK-NEXT: st.param.b32 [func_retval0], %r2;81; CHECK-NEXT: ret;82 %t1 = fdiv afn float %a, 2.083 ret float %t184}85 86define float @rcp_fast(float %a) {87; CHECK-LABEL: rcp_fast(88; CHECK: {89; CHECK-NEXT: .reg .b32 %r<3>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b32 %r1, [rcp_fast_param_0];93; CHECK-NEXT: rcp.approx.f32 %r2, %r1;94; CHECK-NEXT: st.param.b32 [func_retval0], %r2;95; CHECK-NEXT: ret;96 %t1 = fdiv afn float 1.0, %a97 ret float %t198}99 100define float @rcp_fast_ftz(float %a) #0 {101; CHECK-LABEL: rcp_fast_ftz(102; CHECK: {103; CHECK-NEXT: .reg .b32 %r<3>;104; CHECK-EMPTY:105; CHECK-NEXT: // %bb.0:106; CHECK-NEXT: ld.param.b32 %r1, [rcp_fast_ftz_param_0];107; CHECK-NEXT: rcp.approx.ftz.f32 %r2, %r1;108; CHECK-NEXT: st.param.b32 [func_retval0], %r2;109; CHECK-NEXT: ret;110 %t1 = fdiv afn float 1.0, %a111 ret float %t1112}113 114define float @div_fast_vec(float %a, float %b, float %c, float %d) {115; CHECK-LABEL: div_fast_vec(116; CHECK: {117; CHECK-NEXT: .reg .b32 %r<8>;118; CHECK-EMPTY:119; CHECK-NEXT: // %bb.0:120; CHECK-NEXT: ld.param.b32 %r1, [div_fast_vec_param_0];121; CHECK-NEXT: ld.param.b32 %r2, [div_fast_vec_param_1];122; CHECK-NEXT: ld.param.b32 %r3, [div_fast_vec_param_2];123; CHECK-NEXT: ld.param.b32 %r4, [div_fast_vec_param_3];124; CHECK-NEXT: div.approx.f32 %r5, %r2, %r4;125; CHECK-NEXT: div.approx.f32 %r6, %r1, %r3;126; CHECK-NEXT: add.rn.f32 %r7, %r6, %r5;127; CHECK-NEXT: st.param.b32 [func_retval0], %r7;128; CHECK-NEXT: ret;129 %ins_a0 = insertelement <2 x float> poison, float %a, i32 0130 %ins_a1 = insertelement <2 x float> %ins_a0, float %b, i32 1131 %ins_b0 = insertelement <2 x float> poison, float %c, i32 0132 %ins_b1 = insertelement <2 x float> %ins_b0, float %d, i32 1133 %fdiv = fdiv fast <2 x float> %ins_a1, %ins_b1134 %ext0 = extractelement <2 x float> %fdiv, i32 0135 %ext1 = extractelement <2 x float> %fdiv, i32 1136 %fadd = fadd float %ext0, %ext1137 ret float %fadd138}139 140define float @div_fast_vec_ftz(float %a, float %b, float %c, float %d) #0 {141; CHECK-LABEL: div_fast_vec_ftz(142; CHECK: {143; CHECK-NEXT: .reg .b32 %r<8>;144; CHECK-EMPTY:145; CHECK-NEXT: // %bb.0:146; CHECK-NEXT: ld.param.b32 %r1, [div_fast_vec_ftz_param_0];147; CHECK-NEXT: ld.param.b32 %r2, [div_fast_vec_ftz_param_1];148; CHECK-NEXT: ld.param.b32 %r3, [div_fast_vec_ftz_param_2];149; CHECK-NEXT: ld.param.b32 %r4, [div_fast_vec_ftz_param_3];150; CHECK-NEXT: div.approx.ftz.f32 %r5, %r2, %r4;151; CHECK-NEXT: div.approx.ftz.f32 %r6, %r1, %r3;152; CHECK-NEXT: add.rn.ftz.f32 %r7, %r6, %r5;153; CHECK-NEXT: st.param.b32 [func_retval0], %r7;154; CHECK-NEXT: ret;155 %ins_a0 = insertelement <2 x float> poison, float %a, i32 0156 %ins_a1 = insertelement <2 x float> %ins_a0, float %b, i32 1157 %ins_b0 = insertelement <2 x float> poison, float %c, i32 0158 %ins_b1 = insertelement <2 x float> %ins_b0, float %d, i32 1159 %fdiv = fdiv fast <2 x float> %ins_a1, %ins_b1160 %ext0 = extractelement <2 x float> %fdiv, i32 0161 %ext1 = extractelement <2 x float> %fdiv, i32 1162 %fadd = fadd float %ext0, %ext1163 ret float %fadd164}165 166attributes #0 = { "denormal-fp-math-f32" = "preserve-sign" }167