317 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_60 | FileCheck %s3 4 5target triple = "nvptx64-unknown-cuda"6 7define half @frem_f16(half %a, half %b) {8; CHECK-LABEL: frem_f16(9; CHECK: {10; CHECK-NEXT: .reg .pred %p<2>;11; CHECK-NEXT: .reg .b16 %rs<4>;12; CHECK-NEXT: .reg .b32 %r<8>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b16 %rs1, [frem_f16_param_0];16; CHECK-NEXT: ld.param.b16 %rs2, [frem_f16_param_1];17; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;18; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;19; CHECK-NEXT: div.rn.f32 %r3, %r2, %r1;20; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;21; CHECK-NEXT: neg.f32 %r5, %r4;22; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, %r2;23; CHECK-NEXT: testp.infinite.f32 %p1, %r1;24; CHECK-NEXT: selp.f32 %r7, %r2, %r6, %p1;25; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r7;26; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;27; CHECK-NEXT: ret;28 %r = frem half %a, %b29 ret half %r30}31 32define half @frem_f16_fast(half %a, half %b) {33; CHECK-LABEL: frem_f16_fast(34; CHECK: {35; CHECK-NEXT: .reg .b16 %rs<4>;36; CHECK-NEXT: .reg .b32 %r<7>;37; CHECK-EMPTY:38; CHECK-NEXT: // %bb.0:39; CHECK-NEXT: ld.param.b16 %rs1, [frem_f16_fast_param_0];40; CHECK-NEXT: ld.param.b16 %rs2, [frem_f16_fast_param_1];41; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;42; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;43; CHECK-NEXT: div.approx.f32 %r3, %r2, %r1;44; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;45; CHECK-NEXT: neg.f32 %r5, %r4;46; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, %r2;47; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r6;48; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;49; CHECK-NEXT: ret;50 %r = frem afn ninf half %a, %b51 ret half %r52}53 54define float @frem_f32(float %a, float %b) {55; CHECK-LABEL: frem_f32(56; CHECK: {57; CHECK-NEXT: .reg .pred %p<2>;58; CHECK-NEXT: .reg .b32 %r<8>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0:61; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_param_0];62; CHECK-NEXT: ld.param.b32 %r2, [frem_f32_param_1];63; CHECK-NEXT: div.rn.f32 %r3, %r1, %r2;64; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;65; CHECK-NEXT: neg.f32 %r5, %r4;66; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;67; CHECK-NEXT: testp.infinite.f32 %p1, %r2;68; CHECK-NEXT: selp.f32 %r7, %r1, %r6, %p1;69; CHECK-NEXT: st.param.b32 [func_retval0], %r7;70; CHECK-NEXT: ret;71 %r = frem float %a, %b72 ret float %r73}74 75define float @frem_f32_fast(float %a, float %b) {76; CHECK-LABEL: frem_f32_fast(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<7>;79; CHECK-EMPTY:80; CHECK-NEXT: // %bb.0:81; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_fast_param_0];82; CHECK-NEXT: ld.param.b32 %r2, [frem_f32_fast_param_1];83; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;84; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;85; CHECK-NEXT: neg.f32 %r5, %r4;86; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;87; CHECK-NEXT: st.param.b32 [func_retval0], %r6;88; CHECK-NEXT: ret;89 %r = frem afn ninf float %a, %b90 ret float %r91}92 93define double @frem_f64(double %a, double %b) {94; CHECK-LABEL: frem_f64(95; CHECK: {96; CHECK-NEXT: .reg .pred %p<2>;97; CHECK-NEXT: .reg .b64 %rd<8>;98; CHECK-EMPTY:99; CHECK-NEXT: // %bb.0:100; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_param_0];101; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_param_1];102; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;103; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;104; CHECK-NEXT: neg.f64 %rd5, %rd4;105; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;106; CHECK-NEXT: testp.infinite.f64 %p1, %rd2;107; CHECK-NEXT: selp.f64 %rd7, %rd1, %rd6, %p1;108; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;109; CHECK-NEXT: ret;110 %r = frem double %a, %b111 ret double %r112}113 114define double @frem_f64_fast(double %a, double %b) {115; CHECK-LABEL: frem_f64_fast(116; CHECK: {117; CHECK-NEXT: .reg .b64 %rd<7>;118; CHECK-EMPTY:119; CHECK-NEXT: // %bb.0:120; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_fast_param_0];121; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_fast_param_1];122; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;123; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;124; CHECK-NEXT: neg.f64 %rd5, %rd4;125; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;126; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;127; CHECK-NEXT: ret;128 %r = frem afn ninf double %a, %b129 ret double %r130}131 132define half @frem_f16_ninf(half %a, half %b) {133; CHECK-LABEL: frem_f16_ninf(134; CHECK: {135; CHECK-NEXT: .reg .b16 %rs<4>;136; CHECK-NEXT: .reg .b32 %r<7>;137; CHECK-EMPTY:138; CHECK-NEXT: // %bb.0:139; CHECK-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_param_0];140; CHECK-NEXT: ld.param.b16 %rs2, [frem_f16_ninf_param_1];141; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;142; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;143; CHECK-NEXT: div.rn.f32 %r3, %r2, %r1;144; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;145; CHECK-NEXT: neg.f32 %r5, %r4;146; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, %r2;147; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r6;148; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;149; CHECK-NEXT: ret;150 %r = frem ninf half %a, %b151 ret half %r152}153 154define half @frem_f16_ninf_fast(half %a, half %b) {155; CHECK-LABEL: frem_f16_ninf_fast(156; CHECK: {157; CHECK-NEXT: .reg .b16 %rs<4>;158; CHECK-NEXT: .reg .b32 %r<7>;159; CHECK-EMPTY:160; CHECK-NEXT: // %bb.0:161; CHECK-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_fast_param_0];162; CHECK-NEXT: ld.param.b16 %rs2, [frem_f16_ninf_fast_param_1];163; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;164; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;165; CHECK-NEXT: div.approx.f32 %r3, %r2, %r1;166; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;167; CHECK-NEXT: neg.f32 %r5, %r4;168; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, %r2;169; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r6;170; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;171; CHECK-NEXT: ret;172 %r = frem afn ninf half %a, %b173 ret half %r174}175 176define float @frem_f32_ninf(float %a, float %b) {177; CHECK-LABEL: frem_f32_ninf(178; CHECK: {179; CHECK-NEXT: .reg .b32 %r<7>;180; CHECK-EMPTY:181; CHECK-NEXT: // %bb.0:182; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_ninf_param_0];183; CHECK-NEXT: ld.param.b32 %r2, [frem_f32_ninf_param_1];184; CHECK-NEXT: div.rn.f32 %r3, %r1, %r2;185; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;186; CHECK-NEXT: neg.f32 %r5, %r4;187; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;188; CHECK-NEXT: st.param.b32 [func_retval0], %r6;189; CHECK-NEXT: ret;190 %r = frem ninf float %a, %b191 ret float %r192}193 194define float @frem_f32_ninf_fast(float %a, float %b) {195; CHECK-LABEL: frem_f32_ninf_fast(196; CHECK: {197; CHECK-NEXT: .reg .b32 %r<7>;198; CHECK-EMPTY:199; CHECK-NEXT: // %bb.0:200; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_ninf_fast_param_0];201; CHECK-NEXT: ld.param.b32 %r2, [frem_f32_ninf_fast_param_1];202; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2;203; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;204; CHECK-NEXT: neg.f32 %r5, %r4;205; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r2, %r1;206; CHECK-NEXT: st.param.b32 [func_retval0], %r6;207; CHECK-NEXT: ret;208 %r = frem afn ninf float %a, %b209 ret float %r210}211 212define double @frem_f64_ninf(double %a, double %b) {213; CHECK-LABEL: frem_f64_ninf(214; CHECK: {215; CHECK-NEXT: .reg .b64 %rd<7>;216; CHECK-EMPTY:217; CHECK-NEXT: // %bb.0:218; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_ninf_param_0];219; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_ninf_param_1];220; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;221; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;222; CHECK-NEXT: neg.f64 %rd5, %rd4;223; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;224; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;225; CHECK-NEXT: ret;226 %r = frem ninf double %a, %b227 ret double %r228}229 230define double @frem_f64_ninf_fast(double %a, double %b) {231; CHECK-LABEL: frem_f64_ninf_fast(232; CHECK: {233; CHECK-NEXT: .reg .b64 %rd<7>;234; CHECK-EMPTY:235; CHECK-NEXT: // %bb.0:236; CHECK-NEXT: ld.param.b64 %rd1, [frem_f64_ninf_fast_param_0];237; CHECK-NEXT: ld.param.b64 %rd2, [frem_f64_ninf_fast_param_1];238; CHECK-NEXT: div.rn.f64 %rd3, %rd1, %rd2;239; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;240; CHECK-NEXT: neg.f64 %rd5, %rd4;241; CHECK-NEXT: fma.rn.f64 %rd6, %rd5, %rd2, %rd1;242; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;243; CHECK-NEXT: ret;244 %r = frem afn ninf double %a, %b245 ret double %r246}247 248define float @frem_f32_imm1_fast(float %a) {249; CHECK-LABEL: frem_f32_imm1_fast(250; CHECK: {251; CHECK-NEXT: .reg .b32 %r<5>;252; CHECK-EMPTY:253; CHECK-NEXT: // %bb.0:254; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_imm1_fast_param_0];255; CHECK-NEXT: mul.rn.f32 %r2, %r1, 0f3E124925;256; CHECK-NEXT: cvt.rzi.f32.f32 %r3, %r2;257; CHECK-NEXT: fma.rn.f32 %r4, %r3, 0fC0E00000, %r1;258; CHECK-NEXT: st.param.b32 [func_retval0], %r4;259; CHECK-NEXT: ret;260 %r = frem arcp float %a, 7.0261 ret float %r262}263define float @frem_f32_imm1_normal(float %a) {264; CHECK-LABEL: frem_f32_imm1_normal(265; CHECK: {266; CHECK-NEXT: .reg .b32 %r<5>;267; CHECK-EMPTY:268; CHECK-NEXT: // %bb.0:269; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_imm1_normal_param_0];270; CHECK-NEXT: div.rn.f32 %r2, %r1, 0f40E00000;271; CHECK-NEXT: cvt.rzi.f32.f32 %r3, %r2;272; CHECK-NEXT: fma.rn.f32 %r4, %r3, 0fC0E00000, %r1;273; CHECK-NEXT: st.param.b32 [func_retval0], %r4;274; CHECK-NEXT: ret;275 %r = frem float %a, 7.0276 ret float %r277}278 279define float @frem_f32_imm2(float %a) {280; CHECK-LABEL: frem_f32_imm2(281; CHECK: {282; CHECK-NEXT: .reg .pred %p<2>;283; CHECK-NEXT: .reg .b32 %r<8>;284; CHECK-EMPTY:285; CHECK-NEXT: // %bb.0:286; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_imm2_param_0];287; CHECK-NEXT: mov.b32 %r2, 0f40E00000;288; CHECK-NEXT: div.rn.f32 %r3, %r2, %r1;289; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;290; CHECK-NEXT: neg.f32 %r5, %r4;291; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, 0f40E00000;292; CHECK-NEXT: testp.infinite.f32 %p1, %r1;293; CHECK-NEXT: selp.f32 %r7, 0f40E00000, %r6, %p1;294; CHECK-NEXT: st.param.b32 [func_retval0], %r7;295; CHECK-NEXT: ret;296 %r = frem float 7.0, %a297 ret float %r298}299 300define float @frem_f32_imm2_fast(float %a) {301; CHECK-LABEL: frem_f32_imm2_fast(302; CHECK: {303; CHECK-NEXT: .reg .b32 %r<7>;304; CHECK-EMPTY:305; CHECK-NEXT: // %bb.0:306; CHECK-NEXT: ld.param.b32 %r1, [frem_f32_imm2_fast_param_0];307; CHECK-NEXT: mov.b32 %r2, 0f40E00000;308; CHECK-NEXT: div.approx.f32 %r3, %r2, %r1;309; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r3;310; CHECK-NEXT: neg.f32 %r5, %r4;311; CHECK-NEXT: fma.rn.f32 %r6, %r5, %r1, 0f40E00000;312; CHECK-NEXT: st.param.b32 [func_retval0], %r6;313; CHECK-NEXT: ret;314 %r = frem afn ninf float 7.0, %a315 ret float %r316}317