1843 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s | FileCheck %s --check-prefixes=CHECK,CHECK-NOF163; RUN: llc < %s -mcpu=sm_80 -mattr +ptx70 | FileCheck %s --check-prefixes=CHECK,CHECK-F164; RUN: llc < %s -mcpu=sm_80 -mattr +ptx70 --nvptx-no-f16-math | FileCheck %s --check-prefixes=CHECK,CHECK-SM80-NOF165; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}6; RUN: %if ptxas-sm_80 %{ llc < %s -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}7; RUN: %if ptxas-sm_80 %{ llc < %s -mcpu=sm_80 --nvptx-no-f16-math | %ptxas-verify -arch=sm_80 %}8 9target triple = "nvptx64-nvidia-cuda"10 11; Checks that llvm intrinsics for math functions are correctly lowered to PTX.12 13declare float @llvm.ceil.f32(float) #014declare double @llvm.ceil.f64(double) #015declare float @llvm.floor.f32(float) #016declare double @llvm.floor.f64(double) #017declare float @llvm.round.f32(float) #018declare double @llvm.round.f64(double) #019declare float @llvm.nearbyint.f32(float) #020declare double @llvm.nearbyint.f64(double) #021declare float @llvm.rint.f32(float) #022declare double @llvm.rint.f64(double) #023declare float @llvm.roundeven.f32(float) #024declare double @llvm.roundeven.f64(double) #025declare float @llvm.trunc.f32(float) #026declare double @llvm.trunc.f64(double) #027declare float @llvm.fabs.f32(float) #028declare double @llvm.fabs.f64(double) #029declare half @llvm.minnum.f16(half, half) #030declare float @llvm.minnum.f32(float, float) #031declare double @llvm.minnum.f64(double, double) #032declare <2 x half> @llvm.minnum.v2f16(<2 x half>, <2 x half>) #033declare half @llvm.minimum.f16(half, half) #034declare float @llvm.minimum.f32(float, float) #035declare double @llvm.minimum.f64(double, double) #036declare <2 x half> @llvm.minimum.v2f16(<2 x half>, <2 x half>) #037declare half @llvm.maxnum.f16(half, half) #038declare float @llvm.maxnum.f32(float, float) #039declare double @llvm.maxnum.f64(double, double) #040declare <2 x half> @llvm.maxnum.v2f16(<2 x half>, <2 x half>) #041declare half @llvm.maximum.f16(half, half) #042declare float @llvm.maximum.f32(float, float) #043declare double @llvm.maximum.f64(double, double) #044declare <2 x half> @llvm.maximum.v2f16(<2 x half>, <2 x half>) #045declare half @llvm.minimumnum.f16(half, half) #046declare float @llvm.minimumnum.f32(float, float) #047declare double @llvm.minimumnum.f64(double, double) #048declare <2 x half> @llvm.minimumnum.v2f16(<2 x half>, <2 x half>) #049declare half @llvm.maximumnum.f16(half, half) #050declare float @llvm.maximumnum.f32(float, float) #051declare double @llvm.maximumnum.f64(double, double) #052declare <2 x half> @llvm.maximumnum.v2f16(<2 x half>, <2 x half>) #053declare float @llvm.fma.f32(float, float, float) #054declare double @llvm.fma.f64(double, double, double) #055 56; ---- ceil ----57 58define float @ceil_float(float %a) {59; CHECK-LABEL: ceil_float(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<3>;62; CHECK-EMPTY:63; CHECK-NEXT: // %bb.0:64; CHECK-NEXT: ld.param.b32 %r1, [ceil_float_param_0];65; CHECK-NEXT: cvt.rpi.f32.f32 %r2, %r1;66; CHECK-NEXT: st.param.b32 [func_retval0], %r2;67; CHECK-NEXT: ret;68 %b = call float @llvm.ceil.f32(float %a)69 ret float %b70}71 72define float @ceil_float_ftz(float %a) #1 {73; CHECK-LABEL: ceil_float_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, [ceil_float_ftz_param_0];79; CHECK-NEXT: cvt.rpi.ftz.f32.f32 %r2, %r1;80; CHECK-NEXT: st.param.b32 [func_retval0], %r2;81; CHECK-NEXT: ret;82 %b = call float @llvm.ceil.f32(float %a)83 ret float %b84}85 86define double @ceil_double(double %a) {87; CHECK-LABEL: ceil_double(88; CHECK: {89; CHECK-NEXT: .reg .b64 %rd<3>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b64 %rd1, [ceil_double_param_0];93; CHECK-NEXT: cvt.rpi.f64.f64 %rd2, %rd1;94; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;95; CHECK-NEXT: ret;96 %b = call double @llvm.ceil.f64(double %a)97 ret double %b98}99 100; ---- floor ----101 102define float @floor_float(float %a) {103; CHECK-LABEL: floor_float(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<3>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [floor_float_param_0];109; CHECK-NEXT: cvt.rmi.f32.f32 %r2, %r1;110; CHECK-NEXT: st.param.b32 [func_retval0], %r2;111; CHECK-NEXT: ret;112 %b = call float @llvm.floor.f32(float %a)113 ret float %b114}115 116define float @floor_float_ftz(float %a) #1 {117; CHECK-LABEL: floor_float_ftz(118; CHECK: {119; CHECK-NEXT: .reg .b32 %r<3>;120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: ld.param.b32 %r1, [floor_float_ftz_param_0];123; CHECK-NEXT: cvt.rmi.ftz.f32.f32 %r2, %r1;124; CHECK-NEXT: st.param.b32 [func_retval0], %r2;125; CHECK-NEXT: ret;126 %b = call float @llvm.floor.f32(float %a)127 ret float %b128}129 130define double @floor_double(double %a) {131; CHECK-LABEL: floor_double(132; CHECK: {133; CHECK-NEXT: .reg .b64 %rd<3>;134; CHECK-EMPTY:135; CHECK-NEXT: // %bb.0:136; CHECK-NEXT: ld.param.b64 %rd1, [floor_double_param_0];137; CHECK-NEXT: cvt.rmi.f64.f64 %rd2, %rd1;138; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;139; CHECK-NEXT: ret;140 %b = call double @llvm.floor.f64(double %a)141 ret double %b142}143 144; ---- round ----145 146define float @round_float(float %a) {147; check the use of sign mask and 0.5 to implement round148; CHECK-LABEL: round_float(149; CHECK: {150; CHECK-NEXT: .reg .pred %p<3>;151; CHECK-NEXT: .reg .b32 %r<10>;152; CHECK-EMPTY:153; CHECK-NEXT: // %bb.0:154; CHECK-NEXT: ld.param.b32 %r1, [round_float_param_0];155; CHECK-NEXT: and.b32 %r2, %r1, -2147483648;156; CHECK-NEXT: or.b32 %r3, %r2, 1056964608;157; CHECK-NEXT: add.rn.f32 %r4, %r1, %r3;158; CHECK-NEXT: cvt.rzi.f32.f32 %r5, %r4;159; CHECK-NEXT: abs.f32 %r6, %r1;160; CHECK-NEXT: setp.gt.f32 %p1, %r6, 0f4B000000;161; CHECK-NEXT: selp.f32 %r7, %r1, %r5, %p1;162; CHECK-NEXT: cvt.rzi.f32.f32 %r8, %r1;163; CHECK-NEXT: setp.lt.f32 %p2, %r6, 0f3F000000;164; CHECK-NEXT: selp.f32 %r9, %r8, %r7, %p2;165; CHECK-NEXT: st.param.b32 [func_retval0], %r9;166; CHECK-NEXT: ret;167 %b = call float @llvm.round.f32(float %a)168 ret float %b169}170 171define float @round_float_ftz(float %a) #1 {172; check the use of sign mask and 0.5 to implement round173; CHECK-LABEL: round_float_ftz(174; CHECK: {175; CHECK-NEXT: .reg .pred %p<3>;176; CHECK-NEXT: .reg .b32 %r<10>;177; CHECK-EMPTY:178; CHECK-NEXT: // %bb.0:179; CHECK-NEXT: ld.param.b32 %r1, [round_float_ftz_param_0];180; CHECK-NEXT: and.b32 %r2, %r1, -2147483648;181; CHECK-NEXT: or.b32 %r3, %r2, 1056964608;182; CHECK-NEXT: add.rn.ftz.f32 %r4, %r1, %r3;183; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r5, %r4;184; CHECK-NEXT: abs.ftz.f32 %r6, %r1;185; CHECK-NEXT: setp.gt.ftz.f32 %p1, %r6, 0f4B000000;186; CHECK-NEXT: selp.f32 %r7, %r1, %r5, %p1;187; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r8, %r1;188; CHECK-NEXT: setp.lt.ftz.f32 %p2, %r6, 0f3F000000;189; CHECK-NEXT: selp.f32 %r9, %r8, %r7, %p2;190; CHECK-NEXT: st.param.b32 [func_retval0], %r9;191; CHECK-NEXT: ret;192 %b = call float @llvm.round.f32(float %a)193 ret float %b194}195 196define double @round_double(double %a) {197; check the use of 0.5 to implement round198; CHECK-LABEL: round_double(199; CHECK: {200; CHECK-NEXT: .reg .pred %p<3>;201; CHECK-NEXT: .reg .b64 %rd<8>;202; CHECK-EMPTY:203; CHECK-NEXT: // %bb.0:204; CHECK-NEXT: ld.param.b64 %rd1, [round_double_param_0];205; CHECK-NEXT: abs.f64 %rd2, %rd1;206; CHECK-NEXT: setp.lt.f64 %p1, %rd2, 0d3FE0000000000000;207; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FE0000000000000;208; CHECK-NEXT: cvt.rzi.f64.f64 %rd4, %rd3;209; CHECK-NEXT: selp.f64 %rd5, 0d0000000000000000, %rd4, %p1;210; CHECK-NEXT: copysign.f64 %rd6, %rd1, %rd5;211; CHECK-NEXT: setp.gt.f64 %p2, %rd2, 0d4330000000000000;212; CHECK-NEXT: selp.f64 %rd7, %rd1, %rd6, %p2;213; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;214; CHECK-NEXT: ret;215 %b = call double @llvm.round.f64(double %a)216 ret double %b217}218 219; ---- nearbyint ----220 221define float @nearbyint_float(float %a) {222; CHECK-LABEL: nearbyint_float(223; CHECK: {224; CHECK-NEXT: .reg .b32 %r<3>;225; CHECK-EMPTY:226; CHECK-NEXT: // %bb.0:227; CHECK-NEXT: ld.param.b32 %r1, [nearbyint_float_param_0];228; CHECK-NEXT: cvt.rni.f32.f32 %r2, %r1;229; CHECK-NEXT: st.param.b32 [func_retval0], %r2;230; CHECK-NEXT: ret;231 %b = call float @llvm.nearbyint.f32(float %a)232 ret float %b233}234 235define float @nearbyint_float_ftz(float %a) #1 {236; CHECK-LABEL: nearbyint_float_ftz(237; CHECK: {238; CHECK-NEXT: .reg .b32 %r<3>;239; CHECK-EMPTY:240; CHECK-NEXT: // %bb.0:241; CHECK-NEXT: ld.param.b32 %r1, [nearbyint_float_ftz_param_0];242; CHECK-NEXT: cvt.rni.ftz.f32.f32 %r2, %r1;243; CHECK-NEXT: st.param.b32 [func_retval0], %r2;244; CHECK-NEXT: ret;245 %b = call float @llvm.nearbyint.f32(float %a)246 ret float %b247}248 249define double @nearbyint_double(double %a) {250; CHECK-LABEL: nearbyint_double(251; CHECK: {252; CHECK-NEXT: .reg .b64 %rd<3>;253; CHECK-EMPTY:254; CHECK-NEXT: // %bb.0:255; CHECK-NEXT: ld.param.b64 %rd1, [nearbyint_double_param_0];256; CHECK-NEXT: cvt.rni.f64.f64 %rd2, %rd1;257; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;258; CHECK-NEXT: ret;259 %b = call double @llvm.nearbyint.f64(double %a)260 ret double %b261}262 263; ---- rint ----264 265define float @rint_float(float %a) {266; CHECK-LABEL: rint_float(267; CHECK: {268; CHECK-NEXT: .reg .b32 %r<3>;269; CHECK-EMPTY:270; CHECK-NEXT: // %bb.0:271; CHECK-NEXT: ld.param.b32 %r1, [rint_float_param_0];272; CHECK-NEXT: cvt.rni.f32.f32 %r2, %r1;273; CHECK-NEXT: st.param.b32 [func_retval0], %r2;274; CHECK-NEXT: ret;275 %b = call float @llvm.rint.f32(float %a)276 ret float %b277}278 279define float @rint_float_ftz(float %a) #1 {280; CHECK-LABEL: rint_float_ftz(281; CHECK: {282; CHECK-NEXT: .reg .b32 %r<3>;283; CHECK-EMPTY:284; CHECK-NEXT: // %bb.0:285; CHECK-NEXT: ld.param.b32 %r1, [rint_float_ftz_param_0];286; CHECK-NEXT: cvt.rni.ftz.f32.f32 %r2, %r1;287; CHECK-NEXT: st.param.b32 [func_retval0], %r2;288; CHECK-NEXT: ret;289 %b = call float @llvm.rint.f32(float %a)290 ret float %b291}292 293define double @rint_double(double %a) {294; CHECK-LABEL: rint_double(295; CHECK: {296; CHECK-NEXT: .reg .b64 %rd<3>;297; CHECK-EMPTY:298; CHECK-NEXT: // %bb.0:299; CHECK-NEXT: ld.param.b64 %rd1, [rint_double_param_0];300; CHECK-NEXT: cvt.rni.f64.f64 %rd2, %rd1;301; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;302; CHECK-NEXT: ret;303 %b = call double @llvm.rint.f64(double %a)304 ret double %b305}306 307; ---- roundeven ----308 309define float @roundeven_float(float %a) {310; CHECK-LABEL: roundeven_float(311; CHECK: {312; CHECK-NEXT: .reg .b32 %r<3>;313; CHECK-EMPTY:314; CHECK-NEXT: // %bb.0:315; CHECK-NEXT: ld.param.b32 %r1, [roundeven_float_param_0];316; CHECK-NEXT: cvt.rni.f32.f32 %r2, %r1;317; CHECK-NEXT: st.param.b32 [func_retval0], %r2;318; CHECK-NEXT: ret;319 %b = call float @llvm.roundeven.f32(float %a)320 ret float %b321}322 323define float @roundeven_float_ftz(float %a) #1 {324; CHECK-LABEL: roundeven_float_ftz(325; CHECK: {326; CHECK-NEXT: .reg .b32 %r<3>;327; CHECK-EMPTY:328; CHECK-NEXT: // %bb.0:329; CHECK-NEXT: ld.param.b32 %r1, [roundeven_float_ftz_param_0];330; CHECK-NEXT: cvt.rni.ftz.f32.f32 %r2, %r1;331; CHECK-NEXT: st.param.b32 [func_retval0], %r2;332; CHECK-NEXT: ret;333 %b = call float @llvm.roundeven.f32(float %a)334 ret float %b335}336 337define double @roundeven_double(double %a) {338; CHECK-LABEL: roundeven_double(339; CHECK: {340; CHECK-NEXT: .reg .b64 %rd<3>;341; CHECK-EMPTY:342; CHECK-NEXT: // %bb.0:343; CHECK-NEXT: ld.param.b64 %rd1, [roundeven_double_param_0];344; CHECK-NEXT: cvt.rni.f64.f64 %rd2, %rd1;345; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;346; CHECK-NEXT: ret;347 %b = call double @llvm.roundeven.f64(double %a)348 ret double %b349}350 351; ---- trunc ----352 353define float @trunc_float(float %a) {354; CHECK-LABEL: trunc_float(355; CHECK: {356; CHECK-NEXT: .reg .b32 %r<3>;357; CHECK-EMPTY:358; CHECK-NEXT: // %bb.0:359; CHECK-NEXT: ld.param.b32 %r1, [trunc_float_param_0];360; CHECK-NEXT: cvt.rzi.f32.f32 %r2, %r1;361; CHECK-NEXT: st.param.b32 [func_retval0], %r2;362; CHECK-NEXT: ret;363 %b = call float @llvm.trunc.f32(float %a)364 ret float %b365}366 367define float @trunc_float_ftz(float %a) #1 {368; CHECK-LABEL: trunc_float_ftz(369; CHECK: {370; CHECK-NEXT: .reg .b32 %r<3>;371; CHECK-EMPTY:372; CHECK-NEXT: // %bb.0:373; CHECK-NEXT: ld.param.b32 %r1, [trunc_float_ftz_param_0];374; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r2, %r1;375; CHECK-NEXT: st.param.b32 [func_retval0], %r2;376; CHECK-NEXT: ret;377 %b = call float @llvm.trunc.f32(float %a)378 ret float %b379}380 381define double @trunc_double(double %a) {382; CHECK-LABEL: trunc_double(383; CHECK: {384; CHECK-NEXT: .reg .b64 %rd<3>;385; CHECK-EMPTY:386; CHECK-NEXT: // %bb.0:387; CHECK-NEXT: ld.param.b64 %rd1, [trunc_double_param_0];388; CHECK-NEXT: cvt.rzi.f64.f64 %rd2, %rd1;389; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;390; CHECK-NEXT: ret;391 %b = call double @llvm.trunc.f64(double %a)392 ret double %b393}394 395; ---- abs ----396 397define float @abs_float(float %a) {398; CHECK-LABEL: abs_float(399; CHECK: {400; CHECK-NEXT: .reg .b32 %r<3>;401; CHECK-EMPTY:402; CHECK-NEXT: // %bb.0:403; CHECK-NEXT: ld.param.b32 %r1, [abs_float_param_0];404; CHECK-NEXT: abs.f32 %r2, %r1;405; CHECK-NEXT: st.param.b32 [func_retval0], %r2;406; CHECK-NEXT: ret;407 %b = call float @llvm.fabs.f32(float %a)408 ret float %b409}410 411define float @abs_float_ftz(float %a) #1 {412; CHECK-LABEL: abs_float_ftz(413; CHECK: {414; CHECK-NEXT: .reg .b32 %r<3>;415; CHECK-EMPTY:416; CHECK-NEXT: // %bb.0:417; CHECK-NEXT: ld.param.b32 %r1, [abs_float_ftz_param_0];418; CHECK-NEXT: abs.ftz.f32 %r2, %r1;419; CHECK-NEXT: st.param.b32 [func_retval0], %r2;420; CHECK-NEXT: ret;421 %b = call float @llvm.fabs.f32(float %a)422 ret float %b423}424 425define double @abs_double(double %a) {426; CHECK-LABEL: abs_double(427; CHECK: {428; CHECK-NEXT: .reg .b64 %rd<3>;429; CHECK-EMPTY:430; CHECK-NEXT: // %bb.0:431; CHECK-NEXT: ld.param.b64 %rd1, [abs_double_param_0];432; CHECK-NEXT: abs.f64 %rd2, %rd1;433; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;434; CHECK-NEXT: ret;435 %b = call double @llvm.fabs.f64(double %a)436 ret double %b437}438 439; ---- minnum ----440 441define half @minnum_half(half %a, half %b) {442; CHECK-NOF16-LABEL: minnum_half(443; CHECK-NOF16: {444; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;445; CHECK-NOF16-NEXT: .reg .b32 %r<4>;446; CHECK-NOF16-EMPTY:447; CHECK-NOF16-NEXT: // %bb.0:448; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];449; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];450; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;451; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;452; CHECK-NOF16-NEXT: min.f32 %r3, %r2, %r1;453; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;454; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;455; CHECK-NOF16-NEXT: ret;456;457; CHECK-F16-LABEL: minnum_half(458; CHECK-F16: {459; CHECK-F16-NEXT: .reg .b16 %rs<4>;460; CHECK-F16-EMPTY:461; CHECK-F16-NEXT: // %bb.0:462; CHECK-F16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];463; CHECK-F16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];464; CHECK-F16-NEXT: min.f16 %rs3, %rs1, %rs2;465; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;466; CHECK-F16-NEXT: ret;467;468; CHECK-SM80-NOF16-LABEL: minnum_half(469; CHECK-SM80-NOF16: {470; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;471; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;472; CHECK-SM80-NOF16-EMPTY:473; CHECK-SM80-NOF16-NEXT: // %bb.0:474; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];475; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];476; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;477; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;478; CHECK-SM80-NOF16-NEXT: min.f32 %r3, %r2, %r1;479; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;480; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;481; CHECK-SM80-NOF16-NEXT: ret;482 %x = call half @llvm.minnum.f16(half %a, half %b)483 ret half %x484}485 486define float @minnum_float(float %a, float %b) {487; CHECK-LABEL: minnum_float(488; CHECK: {489; CHECK-NEXT: .reg .b32 %r<4>;490; CHECK-EMPTY:491; CHECK-NEXT: // %bb.0:492; CHECK-NEXT: ld.param.b32 %r1, [minnum_float_param_0];493; CHECK-NEXT: ld.param.b32 %r2, [minnum_float_param_1];494; CHECK-NEXT: min.f32 %r3, %r1, %r2;495; CHECK-NEXT: st.param.b32 [func_retval0], %r3;496; CHECK-NEXT: ret;497 %x = call float @llvm.minnum.f32(float %a, float %b)498 ret float %x499}500 501define float @minnum_imm1(float %a) {502; CHECK-LABEL: minnum_imm1(503; CHECK: {504; CHECK-NEXT: .reg .b32 %r<3>;505; CHECK-EMPTY:506; CHECK-NEXT: // %bb.0:507; CHECK-NEXT: ld.param.b32 %r1, [minnum_imm1_param_0];508; CHECK-NEXT: min.f32 %r2, %r1, 0f00000000;509; CHECK-NEXT: st.param.b32 [func_retval0], %r2;510; CHECK-NEXT: ret;511 %x = call float @llvm.minnum.f32(float %a, float 0.0)512 ret float %x513}514 515define float @minnum_imm2(float %a) {516; CHECK-LABEL: minnum_imm2(517; CHECK: {518; CHECK-NEXT: .reg .b32 %r<3>;519; CHECK-EMPTY:520; CHECK-NEXT: // %bb.0:521; CHECK-NEXT: ld.param.b32 %r1, [minnum_imm2_param_0];522; CHECK-NEXT: min.f32 %r2, %r1, 0f00000000;523; CHECK-NEXT: st.param.b32 [func_retval0], %r2;524; CHECK-NEXT: ret;525 %x = call float @llvm.minnum.f32(float 0.0, float %a)526 ret float %x527}528 529define float @minnum_float_ftz(float %a, float %b) #1 {530; CHECK-LABEL: minnum_float_ftz(531; CHECK: {532; CHECK-NEXT: .reg .b32 %r<4>;533; CHECK-EMPTY:534; CHECK-NEXT: // %bb.0:535; CHECK-NEXT: ld.param.b32 %r1, [minnum_float_ftz_param_0];536; CHECK-NEXT: ld.param.b32 %r2, [minnum_float_ftz_param_1];537; CHECK-NEXT: min.ftz.f32 %r3, %r1, %r2;538; CHECK-NEXT: st.param.b32 [func_retval0], %r3;539; CHECK-NEXT: ret;540 %x = call float @llvm.minnum.f32(float %a, float %b)541 ret float %x542}543 544define double @minnum_double(double %a, double %b) {545; CHECK-LABEL: minnum_double(546; CHECK: {547; CHECK-NEXT: .reg .b64 %rd<4>;548; CHECK-EMPTY:549; CHECK-NEXT: // %bb.0:550; CHECK-NEXT: ld.param.b64 %rd1, [minnum_double_param_0];551; CHECK-NEXT: ld.param.b64 %rd2, [minnum_double_param_1];552; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;553; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;554; CHECK-NEXT: ret;555 %x = call double @llvm.minnum.f64(double %a, double %b)556 ret double %x557}558 559define <2 x half> @minnum_v2half(<2 x half> %a, <2 x half> %b) {560; CHECK-NOF16-LABEL: minnum_v2half(561; CHECK-NOF16: {562; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;563; CHECK-NOF16-NEXT: .reg .b32 %r<8>;564; CHECK-NOF16-EMPTY:565; CHECK-NOF16-NEXT: // %bb.0:566; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minnum_v2half_param_0];567; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minnum_v2half_param_1];568; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;569; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;570; CHECK-NOF16-NEXT: min.f32 %r3, %r2, %r1;571; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;572; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;573; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;574; CHECK-NOF16-NEXT: min.f32 %r6, %r5, %r4;575; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;576; CHECK-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};577; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;578; CHECK-NOF16-NEXT: ret;579;580; CHECK-F16-LABEL: minnum_v2half(581; CHECK-F16: {582; CHECK-F16-NEXT: .reg .b32 %r<4>;583; CHECK-F16-EMPTY:584; CHECK-F16-NEXT: // %bb.0:585; CHECK-F16-NEXT: ld.param.b32 %r1, [minnum_v2half_param_0];586; CHECK-F16-NEXT: ld.param.b32 %r2, [minnum_v2half_param_1];587; CHECK-F16-NEXT: min.f16x2 %r3, %r1, %r2;588; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;589; CHECK-F16-NEXT: ret;590;591; CHECK-SM80-NOF16-LABEL: minnum_v2half(592; CHECK-SM80-NOF16: {593; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;594; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<8>;595; CHECK-SM80-NOF16-EMPTY:596; CHECK-SM80-NOF16-NEXT: // %bb.0:597; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minnum_v2half_param_0];598; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minnum_v2half_param_1];599; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;600; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;601; CHECK-SM80-NOF16-NEXT: min.f32 %r3, %r2, %r1;602; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;603; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;604; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;605; CHECK-SM80-NOF16-NEXT: min.f32 %r6, %r5, %r4;606; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;607; CHECK-SM80-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};608; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r7;609; CHECK-SM80-NOF16-NEXT: ret;610 %x = call <2 x half> @llvm.minnum.v2f16(<2 x half> %a, <2 x half> %b)611 ret <2 x half> %x612}613 614; ---- minimum ----615 616define half @minimum_half(half %a, half %b) {617; CHECK-NOF16-LABEL: minimum_half(618; CHECK-NOF16: {619; CHECK-NOF16-NEXT: .reg .pred %p<6>;620; CHECK-NOF16-NEXT: .reg .b16 %rs<8>;621; CHECK-NOF16-NEXT: .reg .b32 %r<4>;622; CHECK-NOF16-EMPTY:623; CHECK-NOF16-NEXT: // %bb.0:624; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];625; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [minimum_half_param_1];626; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;627; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;628; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %r2, %r1;629; CHECK-NOF16-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;630; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;631; CHECK-NOF16-NEXT: selp.b16 %rs4, 0x7E00, %rs3, %p2;632; CHECK-NOF16-NEXT: setp.eq.b16 %p3, %rs1, -32768;633; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs1, %rs4, %p3;634; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs2, -32768;635; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs2, %rs5, %p4;636; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;637; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;638; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs6, %rs4, %p5;639; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs7;640; CHECK-NOF16-NEXT: ret;641;642; CHECK-F16-LABEL: minimum_half(643; CHECK-F16: {644; CHECK-F16-NEXT: .reg .b16 %rs<4>;645; CHECK-F16-EMPTY:646; CHECK-F16-NEXT: // %bb.0:647; CHECK-F16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];648; CHECK-F16-NEXT: ld.param.b16 %rs2, [minimum_half_param_1];649; CHECK-F16-NEXT: min.NaN.f16 %rs3, %rs1, %rs2;650; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;651; CHECK-F16-NEXT: ret;652;653; CHECK-SM80-NOF16-LABEL: minimum_half(654; CHECK-SM80-NOF16: {655; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;656; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<8>;657; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;658; CHECK-SM80-NOF16-EMPTY:659; CHECK-SM80-NOF16-NEXT: // %bb.0:660; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];661; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [minimum_half_param_1];662; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;663; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;664; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p1, %r2, %r1;665; CHECK-SM80-NOF16-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;666; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;667; CHECK-SM80-NOF16-NEXT: selp.b16 %rs4, 0x7E00, %rs3, %p2;668; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p3, %rs1, -32768;669; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs1, %rs4, %p3;670; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs2, -32768;671; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs2, %rs5, %p4;672; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;673; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;674; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs6, %rs4, %p5;675; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs7;676; CHECK-SM80-NOF16-NEXT: ret;677 %x = call half @llvm.minimum.f16(half %a, half %b)678 ret half %x679}680 681define float @minimum_float(float %a, float %b) {682; CHECK-NOF16-LABEL: minimum_float(683; CHECK-NOF16: {684; CHECK-NOF16-NEXT: .reg .pred %p<5>;685; CHECK-NOF16-NEXT: .reg .b32 %r<8>;686; CHECK-NOF16-EMPTY:687; CHECK-NOF16-NEXT: // %bb.0:688; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_param_0];689; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_param_1];690; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;691; CHECK-NOF16-NEXT: min.f32 %r3, %r1, %r2;692; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;693; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;694; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;695; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, -2147483648;696; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;697; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %r4, 0f00000000;698; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;699; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;700; CHECK-NOF16-NEXT: ret;701;702; CHECK-F16-LABEL: minimum_float(703; CHECK-F16: {704; CHECK-F16-NEXT: .reg .b32 %r<4>;705; CHECK-F16-EMPTY:706; CHECK-F16-NEXT: // %bb.0:707; CHECK-F16-NEXT: ld.param.b32 %r1, [minimum_float_param_0];708; CHECK-F16-NEXT: ld.param.b32 %r2, [minimum_float_param_1];709; CHECK-F16-NEXT: min.NaN.f32 %r3, %r1, %r2;710; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;711; CHECK-F16-NEXT: ret;712;713; CHECK-SM80-NOF16-LABEL: minimum_float(714; CHECK-SM80-NOF16: {715; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;716; CHECK-SM80-NOF16-EMPTY:717; CHECK-SM80-NOF16-NEXT: // %bb.0:718; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_param_0];719; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_param_1];720; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %r3, %r1, %r2;721; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r3;722; CHECK-SM80-NOF16-NEXT: ret;723 %x = call float @llvm.minimum.f32(float %a, float %b)724 ret float %x725}726 727define float @minimum_imm1(float %a) {728; CHECK-NOF16-LABEL: minimum_imm1(729; CHECK-NOF16: {730; CHECK-NOF16-NEXT: .reg .pred %p<4>;731; CHECK-NOF16-NEXT: .reg .b32 %r<6>;732; CHECK-NOF16-EMPTY:733; CHECK-NOF16-NEXT: // %bb.0:734; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm1_param_0];735; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;736; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;737; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;738; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;739; CHECK-NOF16-NEXT: selp.f32 %r4, %r1, %r3, %p2;740; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %r3, 0f00000000;741; CHECK-NOF16-NEXT: selp.f32 %r5, %r4, %r3, %p3;742; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r5;743; CHECK-NOF16-NEXT: ret;744;745; CHECK-F16-LABEL: minimum_imm1(746; CHECK-F16: {747; CHECK-F16-NEXT: .reg .b32 %r<3>;748; CHECK-F16-EMPTY:749; CHECK-F16-NEXT: // %bb.0:750; CHECK-F16-NEXT: ld.param.b32 %r1, [minimum_imm1_param_0];751; CHECK-F16-NEXT: min.NaN.f32 %r2, %r1, 0f00000000;752; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r2;753; CHECK-F16-NEXT: ret;754;755; CHECK-SM80-NOF16-LABEL: minimum_imm1(756; CHECK-SM80-NOF16: {757; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<3>;758; CHECK-SM80-NOF16-EMPTY:759; CHECK-SM80-NOF16-NEXT: // %bb.0:760; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm1_param_0];761; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %r2, %r1, 0f00000000;762; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r2;763; CHECK-SM80-NOF16-NEXT: ret;764 %x = call float @llvm.minimum.f32(float %a, float 0.0)765 ret float %x766}767 768define float @minimum_imm2(float %a) {769; CHECK-NOF16-LABEL: minimum_imm2(770; CHECK-NOF16: {771; CHECK-NOF16-NEXT: .reg .pred %p<4>;772; CHECK-NOF16-NEXT: .reg .b32 %r<6>;773; CHECK-NOF16-EMPTY:774; CHECK-NOF16-NEXT: // %bb.0:775; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm2_param_0];776; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;777; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;778; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;779; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;780; CHECK-NOF16-NEXT: selp.f32 %r4, %r1, %r3, %p2;781; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %r3, 0f00000000;782; CHECK-NOF16-NEXT: selp.f32 %r5, %r4, %r3, %p3;783; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r5;784; CHECK-NOF16-NEXT: ret;785;786; CHECK-F16-LABEL: minimum_imm2(787; CHECK-F16: {788; CHECK-F16-NEXT: .reg .b32 %r<3>;789; CHECK-F16-EMPTY:790; CHECK-F16-NEXT: // %bb.0:791; CHECK-F16-NEXT: ld.param.b32 %r1, [minimum_imm2_param_0];792; CHECK-F16-NEXT: min.NaN.f32 %r2, %r1, 0f00000000;793; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r2;794; CHECK-F16-NEXT: ret;795;796; CHECK-SM80-NOF16-LABEL: minimum_imm2(797; CHECK-SM80-NOF16: {798; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<3>;799; CHECK-SM80-NOF16-EMPTY:800; CHECK-SM80-NOF16-NEXT: // %bb.0:801; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm2_param_0];802; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %r2, %r1, 0f00000000;803; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r2;804; CHECK-SM80-NOF16-NEXT: ret;805 %x = call float @llvm.minimum.f32(float 0.0, float %a)806 ret float %x807}808 809define float @minimum_float_ftz(float %a, float %b) #1 {810; CHECK-NOF16-LABEL: minimum_float_ftz(811; CHECK-NOF16: {812; CHECK-NOF16-NEXT: .reg .pred %p<5>;813; CHECK-NOF16-NEXT: .reg .b32 %r<8>;814; CHECK-NOF16-EMPTY:815; CHECK-NOF16-NEXT: // %bb.0:816; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_ftz_param_0];817; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_ftz_param_1];818; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;819; CHECK-NOF16-NEXT: min.ftz.f32 %r3, %r1, %r2;820; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;821; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;822; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;823; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, -2147483648;824; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;825; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %r4, 0f00000000;826; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;827; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;828; CHECK-NOF16-NEXT: ret;829;830; CHECK-F16-LABEL: minimum_float_ftz(831; CHECK-F16: {832; CHECK-F16-NEXT: .reg .b32 %r<4>;833; CHECK-F16-EMPTY:834; CHECK-F16-NEXT: // %bb.0:835; CHECK-F16-NEXT: ld.param.b32 %r1, [minimum_float_ftz_param_0];836; CHECK-F16-NEXT: ld.param.b32 %r2, [minimum_float_ftz_param_1];837; CHECK-F16-NEXT: min.ftz.NaN.f32 %r3, %r1, %r2;838; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;839; CHECK-F16-NEXT: ret;840;841; CHECK-SM80-NOF16-LABEL: minimum_float_ftz(842; CHECK-SM80-NOF16: {843; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;844; CHECK-SM80-NOF16-EMPTY:845; CHECK-SM80-NOF16-NEXT: // %bb.0:846; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_ftz_param_0];847; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_ftz_param_1];848; CHECK-SM80-NOF16-NEXT: min.ftz.NaN.f32 %r3, %r1, %r2;849; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r3;850; CHECK-SM80-NOF16-NEXT: ret;851 %x = call float @llvm.minimum.f32(float %a, float %b)852 ret float %x853}854 855define double @minimum_double(double %a, double %b) {856; CHECK-LABEL: minimum_double(857; CHECK: {858; CHECK-NEXT: .reg .pred %p<5>;859; CHECK-NEXT: .reg .b64 %rd<8>;860; CHECK-EMPTY:861; CHECK-NEXT: // %bb.0:862; CHECK-NEXT: ld.param.b64 %rd1, [minimum_double_param_0];863; CHECK-NEXT: ld.param.b64 %rd2, [minimum_double_param_1];864; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;865; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;866; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;867; CHECK-NEXT: setp.eq.b64 %p2, %rd1, -9223372036854775808;868; CHECK-NEXT: selp.f64 %rd5, %rd1, %rd4, %p2;869; CHECK-NEXT: setp.eq.b64 %p3, %rd2, -9223372036854775808;870; CHECK-NEXT: selp.f64 %rd6, %rd2, %rd5, %p3;871; CHECK-NEXT: setp.eq.f64 %p4, %rd4, 0d0000000000000000;872; CHECK-NEXT: selp.f64 %rd7, %rd6, %rd4, %p4;873; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;874; CHECK-NEXT: ret;875 %x = call double @llvm.minimum.f64(double %a, double %b)876 ret double %x877}878 879define <2 x half> @minimum_v2half(<2 x half> %a, <2 x half> %b) {880; CHECK-NOF16-LABEL: minimum_v2half(881; CHECK-NOF16: {882; CHECK-NOF16-NEXT: .reg .pred %p<11>;883; CHECK-NOF16-NEXT: .reg .b16 %rs<15>;884; CHECK-NOF16-NEXT: .reg .b32 %r<7>;885; CHECK-NOF16-EMPTY:886; CHECK-NOF16-NEXT: // %bb.0:887; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimum_v2half_param_0];888; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimum_v2half_param_1];889; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;890; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;891; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %r2, %r1;892; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p1;893; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;894; CHECK-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;895; CHECK-NOF16-NEXT: setp.eq.b16 %p3, %rs2, -32768;896; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs2, %rs6, %p3;897; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs4, -32768;898; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs7, %p4;899; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs6;900; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;901; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p5;902; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;903; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;904; CHECK-NOF16-NEXT: setp.lt.f32 %p6, %r5, %r4;905; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs1, %rs3, %p6;906; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %r5, %r4;907; CHECK-NOF16-NEXT: selp.b16 %rs11, 0x7E00, %rs10, %p7;908; CHECK-NOF16-NEXT: setp.eq.b16 %p8, %rs1, -32768;909; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs1, %rs11, %p8;910; CHECK-NOF16-NEXT: setp.eq.b16 %p9, %rs3, -32768;911; CHECK-NOF16-NEXT: selp.b16 %rs13, %rs3, %rs12, %p9;912; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs11;913; CHECK-NOF16-NEXT: setp.eq.f32 %p10, %r6, 0f00000000;914; CHECK-NOF16-NEXT: selp.b16 %rs14, %rs13, %rs11, %p10;915; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs14, %rs9};916; CHECK-NOF16-NEXT: ret;917;918; CHECK-F16-LABEL: minimum_v2half(919; CHECK-F16: {920; CHECK-F16-NEXT: .reg .b32 %r<4>;921; CHECK-F16-EMPTY:922; CHECK-F16-NEXT: // %bb.0:923; CHECK-F16-NEXT: ld.param.b32 %r1, [minimum_v2half_param_0];924; CHECK-F16-NEXT: ld.param.b32 %r2, [minimum_v2half_param_1];925; CHECK-F16-NEXT: min.NaN.f16x2 %r3, %r1, %r2;926; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;927; CHECK-F16-NEXT: ret;928;929; CHECK-SM80-NOF16-LABEL: minimum_v2half(930; CHECK-SM80-NOF16: {931; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;932; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<15>;933; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<7>;934; CHECK-SM80-NOF16-EMPTY:935; CHECK-SM80-NOF16-NEXT: // %bb.0:936; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimum_v2half_param_0];937; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimum_v2half_param_1];938; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;939; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;940; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p1, %r2, %r1;941; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p1;942; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;943; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;944; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p3, %rs2, -32768;945; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs2, %rs6, %p3;946; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs4, -32768;947; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs7, %p4;948; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs6;949; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;950; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p5;951; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;952; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;953; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p6, %r5, %r4;954; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs1, %rs3, %p6;955; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %r5, %r4;956; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, 0x7E00, %rs10, %p7;957; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p8, %rs1, -32768;958; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs1, %rs11, %p8;959; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p9, %rs3, -32768;960; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, %rs3, %rs12, %p9;961; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r6, %rs11;962; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p10, %r6, 0f00000000;963; CHECK-SM80-NOF16-NEXT: selp.b16 %rs14, %rs13, %rs11, %p10;964; CHECK-SM80-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs14, %rs9};965; CHECK-SM80-NOF16-NEXT: ret;966 %x = call <2 x half> @llvm.minimum.v2f16(<2 x half> %a, <2 x half> %b)967 ret <2 x half> %x968}969 970; ---- maxnum ----971 972define half @maxnum_half(half %a, half %b) {973; CHECK-NOF16-LABEL: maxnum_half(974; CHECK-NOF16: {975; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;976; CHECK-NOF16-NEXT: .reg .b32 %r<4>;977; CHECK-NOF16-EMPTY:978; CHECK-NOF16-NEXT: // %bb.0:979; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];980; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];981; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;982; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;983; CHECK-NOF16-NEXT: max.f32 %r3, %r2, %r1;984; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;985; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;986; CHECK-NOF16-NEXT: ret;987;988; CHECK-F16-LABEL: maxnum_half(989; CHECK-F16: {990; CHECK-F16-NEXT: .reg .b16 %rs<4>;991; CHECK-F16-EMPTY:992; CHECK-F16-NEXT: // %bb.0:993; CHECK-F16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];994; CHECK-F16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];995; CHECK-F16-NEXT: max.f16 %rs3, %rs1, %rs2;996; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;997; CHECK-F16-NEXT: ret;998;999; CHECK-SM80-NOF16-LABEL: maxnum_half(1000; CHECK-SM80-NOF16: {1001; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;1002; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1003; CHECK-SM80-NOF16-EMPTY:1004; CHECK-SM80-NOF16-NEXT: // %bb.0:1005; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];1006; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];1007; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1008; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1009; CHECK-SM80-NOF16-NEXT: max.f32 %r3, %r2, %r1;1010; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1011; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;1012; CHECK-SM80-NOF16-NEXT: ret;1013 %x = call half @llvm.maxnum.f16(half %a, half %b)1014 ret half %x1015}1016 1017define float @maxnum_imm1(float %a) {1018; CHECK-LABEL: maxnum_imm1(1019; CHECK: {1020; CHECK-NEXT: .reg .b32 %r<3>;1021; CHECK-EMPTY:1022; CHECK-NEXT: // %bb.0:1023; CHECK-NEXT: ld.param.b32 %r1, [maxnum_imm1_param_0];1024; CHECK-NEXT: max.f32 %r2, %r1, 0f00000000;1025; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1026; CHECK-NEXT: ret;1027 %x = call float @llvm.maxnum.f32(float %a, float 0.0)1028 ret float %x1029}1030 1031define float @maxnum_imm2(float %a) {1032; CHECK-LABEL: maxnum_imm2(1033; CHECK: {1034; CHECK-NEXT: .reg .b32 %r<3>;1035; CHECK-EMPTY:1036; CHECK-NEXT: // %bb.0:1037; CHECK-NEXT: ld.param.b32 %r1, [maxnum_imm2_param_0];1038; CHECK-NEXT: max.f32 %r2, %r1, 0f00000000;1039; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1040; CHECK-NEXT: ret;1041 %x = call float @llvm.maxnum.f32(float 0.0, float %a)1042 ret float %x1043}1044 1045define float @maxnum_float(float %a, float %b) {1046; CHECK-LABEL: maxnum_float(1047; CHECK: {1048; CHECK-NEXT: .reg .b32 %r<4>;1049; CHECK-EMPTY:1050; CHECK-NEXT: // %bb.0:1051; CHECK-NEXT: ld.param.b32 %r1, [maxnum_float_param_0];1052; CHECK-NEXT: ld.param.b32 %r2, [maxnum_float_param_1];1053; CHECK-NEXT: max.f32 %r3, %r1, %r2;1054; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1055; CHECK-NEXT: ret;1056 %x = call float @llvm.maxnum.f32(float %a, float %b)1057 ret float %x1058}1059 1060define float @maxnum_float_ftz(float %a, float %b) #1 {1061; CHECK-LABEL: maxnum_float_ftz(1062; CHECK: {1063; CHECK-NEXT: .reg .b32 %r<4>;1064; CHECK-EMPTY:1065; CHECK-NEXT: // %bb.0:1066; CHECK-NEXT: ld.param.b32 %r1, [maxnum_float_ftz_param_0];1067; CHECK-NEXT: ld.param.b32 %r2, [maxnum_float_ftz_param_1];1068; CHECK-NEXT: max.ftz.f32 %r3, %r1, %r2;1069; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1070; CHECK-NEXT: ret;1071 %x = call float @llvm.maxnum.f32(float %a, float %b)1072 ret float %x1073}1074 1075define double @maxnum_double(double %a, double %b) {1076; CHECK-LABEL: maxnum_double(1077; CHECK: {1078; CHECK-NEXT: .reg .b64 %rd<4>;1079; CHECK-EMPTY:1080; CHECK-NEXT: // %bb.0:1081; CHECK-NEXT: ld.param.b64 %rd1, [maxnum_double_param_0];1082; CHECK-NEXT: ld.param.b64 %rd2, [maxnum_double_param_1];1083; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;1084; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;1085; CHECK-NEXT: ret;1086 %x = call double @llvm.maxnum.f64(double %a, double %b)1087 ret double %x1088}1089 1090define <2 x half> @maxnum_v2half(<2 x half> %a, <2 x half> %b) {1091; CHECK-NOF16-LABEL: maxnum_v2half(1092; CHECK-NOF16: {1093; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1094; CHECK-NOF16-NEXT: .reg .b32 %r<8>;1095; CHECK-NOF16-EMPTY:1096; CHECK-NOF16-NEXT: // %bb.0:1097; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maxnum_v2half_param_0];1098; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maxnum_v2half_param_1];1099; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1100; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1101; CHECK-NOF16-NEXT: max.f32 %r3, %r2, %r1;1102; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1103; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1104; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1105; CHECK-NOF16-NEXT: max.f32 %r6, %r5, %r4;1106; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1107; CHECK-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1108; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1109; CHECK-NOF16-NEXT: ret;1110;1111; CHECK-F16-LABEL: maxnum_v2half(1112; CHECK-F16: {1113; CHECK-F16-NEXT: .reg .b32 %r<4>;1114; CHECK-F16-EMPTY:1115; CHECK-F16-NEXT: // %bb.0:1116; CHECK-F16-NEXT: ld.param.b32 %r1, [maxnum_v2half_param_0];1117; CHECK-F16-NEXT: ld.param.b32 %r2, [maxnum_v2half_param_1];1118; CHECK-F16-NEXT: max.f16x2 %r3, %r1, %r2;1119; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1120; CHECK-F16-NEXT: ret;1121;1122; CHECK-SM80-NOF16-LABEL: maxnum_v2half(1123; CHECK-SM80-NOF16: {1124; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;1125; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<8>;1126; CHECK-SM80-NOF16-EMPTY:1127; CHECK-SM80-NOF16-NEXT: // %bb.0:1128; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maxnum_v2half_param_0];1129; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maxnum_v2half_param_1];1130; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1131; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1132; CHECK-SM80-NOF16-NEXT: max.f32 %r3, %r2, %r1;1133; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1134; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1135; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1136; CHECK-SM80-NOF16-NEXT: max.f32 %r6, %r5, %r4;1137; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1138; CHECK-SM80-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1139; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1140; CHECK-SM80-NOF16-NEXT: ret;1141 %x = call <2 x half> @llvm.maxnum.v2f16(<2 x half> %a, <2 x half> %b)1142 ret <2 x half> %x1143}1144 1145; ---- maximum ----1146 1147define half @maximum_half(half %a, half %b) {1148; CHECK-NOF16-LABEL: maximum_half(1149; CHECK-NOF16: {1150; CHECK-NOF16-NEXT: .reg .pred %p<6>;1151; CHECK-NOF16-NEXT: .reg .b16 %rs<8>;1152; CHECK-NOF16-NEXT: .reg .b32 %r<4>;1153; CHECK-NOF16-EMPTY:1154; CHECK-NOF16-NEXT: // %bb.0:1155; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];1156; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [maximum_half_param_1];1157; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1158; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1159; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %r2, %r1;1160; CHECK-NOF16-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;1161; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;1162; CHECK-NOF16-NEXT: selp.b16 %rs4, 0x7E00, %rs3, %p2;1163; CHECK-NOF16-NEXT: setp.eq.b16 %p3, %rs1, 0;1164; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs1, %rs4, %p3;1165; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs2, 0;1166; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs2, %rs5, %p4;1167; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;1168; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;1169; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs6, %rs4, %p5;1170; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs7;1171; CHECK-NOF16-NEXT: ret;1172;1173; CHECK-F16-LABEL: maximum_half(1174; CHECK-F16: {1175; CHECK-F16-NEXT: .reg .b16 %rs<4>;1176; CHECK-F16-EMPTY:1177; CHECK-F16-NEXT: // %bb.0:1178; CHECK-F16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];1179; CHECK-F16-NEXT: ld.param.b16 %rs2, [maximum_half_param_1];1180; CHECK-F16-NEXT: max.NaN.f16 %rs3, %rs1, %rs2;1181; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;1182; CHECK-F16-NEXT: ret;1183;1184; CHECK-SM80-NOF16-LABEL: maximum_half(1185; CHECK-SM80-NOF16: {1186; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;1187; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<8>;1188; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1189; CHECK-SM80-NOF16-EMPTY:1190; CHECK-SM80-NOF16-NEXT: // %bb.0:1191; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];1192; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [maximum_half_param_1];1193; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1194; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1195; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p1, %r2, %r1;1196; CHECK-SM80-NOF16-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;1197; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;1198; CHECK-SM80-NOF16-NEXT: selp.b16 %rs4, 0x7E00, %rs3, %p2;1199; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p3, %rs1, 0;1200; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs1, %rs4, %p3;1201; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs2, 0;1202; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs2, %rs5, %p4;1203; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;1204; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;1205; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs6, %rs4, %p5;1206; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs7;1207; CHECK-SM80-NOF16-NEXT: ret;1208 %x = call half @llvm.maximum.f16(half %a, half %b)1209 ret half %x1210}1211 1212define float @maximum_imm1(float %a) {1213; CHECK-NOF16-LABEL: maximum_imm1(1214; CHECK-NOF16: {1215; CHECK-NOF16-NEXT: .reg .pred %p<3>;1216; CHECK-NOF16-NEXT: .reg .b32 %r<5>;1217; CHECK-NOF16-EMPTY:1218; CHECK-NOF16-NEXT: // %bb.0:1219; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm1_param_0];1220; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;1221; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;1222; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;1223; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r3, 0f00000000;1224; CHECK-NOF16-NEXT: selp.f32 %r4, 0f00000000, %r3, %p2;1225; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;1226; CHECK-NOF16-NEXT: ret;1227;1228; CHECK-F16-LABEL: maximum_imm1(1229; CHECK-F16: {1230; CHECK-F16-NEXT: .reg .b32 %r<3>;1231; CHECK-F16-EMPTY:1232; CHECK-F16-NEXT: // %bb.0:1233; CHECK-F16-NEXT: ld.param.b32 %r1, [maximum_imm1_param_0];1234; CHECK-F16-NEXT: max.NaN.f32 %r2, %r1, 0f00000000;1235; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r2;1236; CHECK-F16-NEXT: ret;1237;1238; CHECK-SM80-NOF16-LABEL: maximum_imm1(1239; CHECK-SM80-NOF16: {1240; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<3>;1241; CHECK-SM80-NOF16-EMPTY:1242; CHECK-SM80-NOF16-NEXT: // %bb.0:1243; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm1_param_0];1244; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %r2, %r1, 0f00000000;1245; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r2;1246; CHECK-SM80-NOF16-NEXT: ret;1247 %x = call float @llvm.maximum.f32(float %a, float 0.0)1248 ret float %x1249}1250 1251define float @maximum_imm2(float %a) {1252; CHECK-NOF16-LABEL: maximum_imm2(1253; CHECK-NOF16: {1254; CHECK-NOF16-NEXT: .reg .pred %p<3>;1255; CHECK-NOF16-NEXT: .reg .b32 %r<5>;1256; CHECK-NOF16-EMPTY:1257; CHECK-NOF16-NEXT: // %bb.0:1258; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm2_param_0];1259; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;1260; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;1261; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;1262; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r3, 0f00000000;1263; CHECK-NOF16-NEXT: selp.f32 %r4, 0f00000000, %r3, %p2;1264; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;1265; CHECK-NOF16-NEXT: ret;1266;1267; CHECK-F16-LABEL: maximum_imm2(1268; CHECK-F16: {1269; CHECK-F16-NEXT: .reg .b32 %r<3>;1270; CHECK-F16-EMPTY:1271; CHECK-F16-NEXT: // %bb.0:1272; CHECK-F16-NEXT: ld.param.b32 %r1, [maximum_imm2_param_0];1273; CHECK-F16-NEXT: max.NaN.f32 %r2, %r1, 0f00000000;1274; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r2;1275; CHECK-F16-NEXT: ret;1276;1277; CHECK-SM80-NOF16-LABEL: maximum_imm2(1278; CHECK-SM80-NOF16: {1279; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<3>;1280; CHECK-SM80-NOF16-EMPTY:1281; CHECK-SM80-NOF16-NEXT: // %bb.0:1282; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm2_param_0];1283; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %r2, %r1, 0f00000000;1284; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r2;1285; CHECK-SM80-NOF16-NEXT: ret;1286 %x = call float @llvm.maximum.f32(float 0.0, float %a)1287 ret float %x1288}1289 1290define float @maximum_float(float %a, float %b) {1291; CHECK-NOF16-LABEL: maximum_float(1292; CHECK-NOF16: {1293; CHECK-NOF16-NEXT: .reg .pred %p<5>;1294; CHECK-NOF16-NEXT: .reg .b32 %r<8>;1295; CHECK-NOF16-EMPTY:1296; CHECK-NOF16-NEXT: // %bb.0:1297; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_param_0];1298; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_param_1];1299; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;1300; CHECK-NOF16-NEXT: max.f32 %r3, %r1, %r2;1301; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;1302; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, 0;1303; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;1304; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, 0;1305; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;1306; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %r4, 0f00000000;1307; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;1308; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1309; CHECK-NOF16-NEXT: ret;1310;1311; CHECK-F16-LABEL: maximum_float(1312; CHECK-F16: {1313; CHECK-F16-NEXT: .reg .b32 %r<4>;1314; CHECK-F16-EMPTY:1315; CHECK-F16-NEXT: // %bb.0:1316; CHECK-F16-NEXT: ld.param.b32 %r1, [maximum_float_param_0];1317; CHECK-F16-NEXT: ld.param.b32 %r2, [maximum_float_param_1];1318; CHECK-F16-NEXT: max.NaN.f32 %r3, %r1, %r2;1319; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1320; CHECK-F16-NEXT: ret;1321;1322; CHECK-SM80-NOF16-LABEL: maximum_float(1323; CHECK-SM80-NOF16: {1324; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1325; CHECK-SM80-NOF16-EMPTY:1326; CHECK-SM80-NOF16-NEXT: // %bb.0:1327; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_param_0];1328; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_param_1];1329; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %r3, %r1, %r2;1330; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r3;1331; CHECK-SM80-NOF16-NEXT: ret;1332 %x = call float @llvm.maximum.f32(float %a, float %b)1333 ret float %x1334}1335 1336define float @maximum_float_ftz(float %a, float %b) #1 {1337; CHECK-NOF16-LABEL: maximum_float_ftz(1338; CHECK-NOF16: {1339; CHECK-NOF16-NEXT: .reg .pred %p<5>;1340; CHECK-NOF16-NEXT: .reg .b32 %r<8>;1341; CHECK-NOF16-EMPTY:1342; CHECK-NOF16-NEXT: // %bb.0:1343; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_ftz_param_0];1344; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_ftz_param_1];1345; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;1346; CHECK-NOF16-NEXT: max.ftz.f32 %r3, %r1, %r2;1347; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;1348; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, 0;1349; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;1350; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, 0;1351; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;1352; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %r4, 0f00000000;1353; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;1354; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1355; CHECK-NOF16-NEXT: ret;1356;1357; CHECK-F16-LABEL: maximum_float_ftz(1358; CHECK-F16: {1359; CHECK-F16-NEXT: .reg .b32 %r<4>;1360; CHECK-F16-EMPTY:1361; CHECK-F16-NEXT: // %bb.0:1362; CHECK-F16-NEXT: ld.param.b32 %r1, [maximum_float_ftz_param_0];1363; CHECK-F16-NEXT: ld.param.b32 %r2, [maximum_float_ftz_param_1];1364; CHECK-F16-NEXT: max.ftz.NaN.f32 %r3, %r1, %r2;1365; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1366; CHECK-F16-NEXT: ret;1367;1368; CHECK-SM80-NOF16-LABEL: maximum_float_ftz(1369; CHECK-SM80-NOF16: {1370; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1371; CHECK-SM80-NOF16-EMPTY:1372; CHECK-SM80-NOF16-NEXT: // %bb.0:1373; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_ftz_param_0];1374; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_ftz_param_1];1375; CHECK-SM80-NOF16-NEXT: max.ftz.NaN.f32 %r3, %r1, %r2;1376; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r3;1377; CHECK-SM80-NOF16-NEXT: ret;1378 %x = call float @llvm.maximum.f32(float %a, float %b)1379 ret float %x1380}1381 1382define double @maximum_double(double %a, double %b) {1383; CHECK-LABEL: maximum_double(1384; CHECK: {1385; CHECK-NEXT: .reg .pred %p<5>;1386; CHECK-NEXT: .reg .b64 %rd<8>;1387; CHECK-EMPTY:1388; CHECK-NEXT: // %bb.0:1389; CHECK-NEXT: ld.param.b64 %rd1, [maximum_double_param_0];1390; CHECK-NEXT: ld.param.b64 %rd2, [maximum_double_param_1];1391; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;1392; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;1393; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;1394; CHECK-NEXT: setp.eq.b64 %p2, %rd1, 0;1395; CHECK-NEXT: selp.f64 %rd5, %rd1, %rd4, %p2;1396; CHECK-NEXT: setp.eq.b64 %p3, %rd2, 0;1397; CHECK-NEXT: selp.f64 %rd6, %rd2, %rd5, %p3;1398; CHECK-NEXT: setp.eq.f64 %p4, %rd4, 0d0000000000000000;1399; CHECK-NEXT: selp.f64 %rd7, %rd6, %rd4, %p4;1400; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;1401; CHECK-NEXT: ret;1402 %x = call double @llvm.maximum.f64(double %a, double %b)1403 ret double %x1404}1405 1406define <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) {1407; CHECK-NOF16-LABEL: maximum_v2half(1408; CHECK-NOF16: {1409; CHECK-NOF16-NEXT: .reg .pred %p<11>;1410; CHECK-NOF16-NEXT: .reg .b16 %rs<15>;1411; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1412; CHECK-NOF16-EMPTY:1413; CHECK-NOF16-NEXT: // %bb.0:1414; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximum_v2half_param_0];1415; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximum_v2half_param_1];1416; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1417; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1418; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %r2, %r1;1419; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p1;1420; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;1421; CHECK-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;1422; CHECK-NOF16-NEXT: setp.eq.b16 %p3, %rs2, 0;1423; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs2, %rs6, %p3;1424; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs4, 0;1425; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs7, %p4;1426; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs6;1427; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;1428; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p5;1429; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1430; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1431; CHECK-NOF16-NEXT: setp.gt.f32 %p6, %r5, %r4;1432; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs1, %rs3, %p6;1433; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %r5, %r4;1434; CHECK-NOF16-NEXT: selp.b16 %rs11, 0x7E00, %rs10, %p7;1435; CHECK-NOF16-NEXT: setp.eq.b16 %p8, %rs1, 0;1436; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs1, %rs11, %p8;1437; CHECK-NOF16-NEXT: setp.eq.b16 %p9, %rs3, 0;1438; CHECK-NOF16-NEXT: selp.b16 %rs13, %rs3, %rs12, %p9;1439; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs11;1440; CHECK-NOF16-NEXT: setp.eq.f32 %p10, %r6, 0f00000000;1441; CHECK-NOF16-NEXT: selp.b16 %rs14, %rs13, %rs11, %p10;1442; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs14, %rs9};1443; CHECK-NOF16-NEXT: ret;1444;1445; CHECK-F16-LABEL: maximum_v2half(1446; CHECK-F16: {1447; CHECK-F16-NEXT: .reg .b32 %r<4>;1448; CHECK-F16-EMPTY:1449; CHECK-F16-NEXT: // %bb.0:1450; CHECK-F16-NEXT: ld.param.b32 %r1, [maximum_v2half_param_0];1451; CHECK-F16-NEXT: ld.param.b32 %r2, [maximum_v2half_param_1];1452; CHECK-F16-NEXT: max.NaN.f16x2 %r3, %r1, %r2;1453; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1454; CHECK-F16-NEXT: ret;1455;1456; CHECK-SM80-NOF16-LABEL: maximum_v2half(1457; CHECK-SM80-NOF16: {1458; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;1459; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<15>;1460; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<7>;1461; CHECK-SM80-NOF16-EMPTY:1462; CHECK-SM80-NOF16-NEXT: // %bb.0:1463; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximum_v2half_param_0];1464; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximum_v2half_param_1];1465; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1466; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1467; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p1, %r2, %r1;1468; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p1;1469; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r2, %r1;1470; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;1471; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p3, %rs2, 0;1472; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs2, %rs6, %p3;1473; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs4, 0;1474; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs7, %p4;1475; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs6;1476; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %r3, 0f00000000;1477; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p5;1478; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1479; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1480; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p6, %r5, %r4;1481; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs1, %rs3, %p6;1482; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %r5, %r4;1483; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, 0x7E00, %rs10, %p7;1484; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p8, %rs1, 0;1485; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs1, %rs11, %p8;1486; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p9, %rs3, 0;1487; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, %rs3, %rs12, %p9;1488; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r6, %rs11;1489; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p10, %r6, 0f00000000;1490; CHECK-SM80-NOF16-NEXT: selp.b16 %rs14, %rs13, %rs11, %p10;1491; CHECK-SM80-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs14, %rs9};1492; CHECK-SM80-NOF16-NEXT: ret;1493 %x = call <2 x half> @llvm.maximum.v2f16(<2 x half> %a, <2 x half> %b)1494 ret <2 x half> %x1495}1496 1497; ---- minimumnum ----1498 1499define half @minimumnum_half(half %a, half %b) {1500; CHECK-NOF16-LABEL: minimumnum_half(1501; CHECK-NOF16: {1502; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;1503; CHECK-NOF16-NEXT: .reg .b32 %r<4>;1504; CHECK-NOF16-EMPTY:1505; CHECK-NOF16-NEXT: // %bb.0:1506; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];1507; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];1508; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1509; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1510; CHECK-NOF16-NEXT: min.f32 %r3, %r2, %r1;1511; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1512; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;1513; CHECK-NOF16-NEXT: ret;1514;1515; CHECK-F16-LABEL: minimumnum_half(1516; CHECK-F16: {1517; CHECK-F16-NEXT: .reg .b16 %rs<4>;1518; CHECK-F16-EMPTY:1519; CHECK-F16-NEXT: // %bb.0:1520; CHECK-F16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];1521; CHECK-F16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];1522; CHECK-F16-NEXT: min.f16 %rs3, %rs1, %rs2;1523; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;1524; CHECK-F16-NEXT: ret;1525;1526; CHECK-SM80-NOF16-LABEL: minimumnum_half(1527; CHECK-SM80-NOF16: {1528; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;1529; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1530; CHECK-SM80-NOF16-EMPTY:1531; CHECK-SM80-NOF16-NEXT: // %bb.0:1532; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];1533; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];1534; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1535; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1536; CHECK-SM80-NOF16-NEXT: min.f32 %r3, %r2, %r1;1537; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1538; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;1539; CHECK-SM80-NOF16-NEXT: ret;1540 %x = call half @llvm.minimumnum.f16(half %a, half %b)1541 ret half %x1542}1543 1544define float @minimumnum_float(float %a, float %b) {1545; CHECK-LABEL: minimumnum_float(1546; CHECK: {1547; CHECK-NEXT: .reg .b32 %r<4>;1548; CHECK-EMPTY:1549; CHECK-NEXT: // %bb.0:1550; CHECK-NEXT: ld.param.b32 %r1, [minimumnum_float_param_0];1551; CHECK-NEXT: ld.param.b32 %r2, [minimumnum_float_param_1];1552; CHECK-NEXT: min.f32 %r3, %r1, %r2;1553; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1554; CHECK-NEXT: ret;1555 %x = call float @llvm.minimumnum.f32(float %a, float %b)1556 ret float %x1557}1558 1559define float @minimumnum_float_ftz(float %a, float %b) #1 {1560; CHECK-LABEL: minimumnum_float_ftz(1561; CHECK: {1562; CHECK-NEXT: .reg .b32 %r<4>;1563; CHECK-EMPTY:1564; CHECK-NEXT: // %bb.0:1565; CHECK-NEXT: ld.param.b32 %r1, [minimumnum_float_ftz_param_0];1566; CHECK-NEXT: ld.param.b32 %r2, [minimumnum_float_ftz_param_1];1567; CHECK-NEXT: min.ftz.f32 %r3, %r1, %r2;1568; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1569; CHECK-NEXT: ret;1570 %x = call float @llvm.minimumnum.f32(float %a, float %b)1571 ret float %x1572}1573 1574define double @minimumnum_double(double %a, double %b) {1575; CHECK-LABEL: minimumnum_double(1576; CHECK: {1577; CHECK-NEXT: .reg .b64 %rd<4>;1578; CHECK-EMPTY:1579; CHECK-NEXT: // %bb.0:1580; CHECK-NEXT: ld.param.b64 %rd1, [minimumnum_double_param_0];1581; CHECK-NEXT: ld.param.b64 %rd2, [minimumnum_double_param_1];1582; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;1583; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;1584; CHECK-NEXT: ret;1585 %x = call double @llvm.minimumnum.f64(double %a, double %b)1586 ret double %x1587}1588 1589define <2 x half> @minimumnum_v2half(<2 x half> %a, <2 x half> %b) {1590; CHECK-NOF16-LABEL: minimumnum_v2half(1591; CHECK-NOF16: {1592; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1593; CHECK-NOF16-NEXT: .reg .b32 %r<8>;1594; CHECK-NOF16-EMPTY:1595; CHECK-NOF16-NEXT: // %bb.0:1596; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimumnum_v2half_param_0];1597; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimumnum_v2half_param_1];1598; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1599; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1600; CHECK-NOF16-NEXT: min.f32 %r3, %r2, %r1;1601; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1602; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1603; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1604; CHECK-NOF16-NEXT: min.f32 %r6, %r5, %r4;1605; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1606; CHECK-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1607; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1608; CHECK-NOF16-NEXT: ret;1609;1610; CHECK-F16-LABEL: minimumnum_v2half(1611; CHECK-F16: {1612; CHECK-F16-NEXT: .reg .b32 %r<4>;1613; CHECK-F16-EMPTY:1614; CHECK-F16-NEXT: // %bb.0:1615; CHECK-F16-NEXT: ld.param.b32 %r1, [minimumnum_v2half_param_0];1616; CHECK-F16-NEXT: ld.param.b32 %r2, [minimumnum_v2half_param_1];1617; CHECK-F16-NEXT: min.f16x2 %r3, %r1, %r2;1618; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1619; CHECK-F16-NEXT: ret;1620;1621; CHECK-SM80-NOF16-LABEL: minimumnum_v2half(1622; CHECK-SM80-NOF16: {1623; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;1624; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<8>;1625; CHECK-SM80-NOF16-EMPTY:1626; CHECK-SM80-NOF16-NEXT: // %bb.0:1627; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimumnum_v2half_param_0];1628; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimumnum_v2half_param_1];1629; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1630; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1631; CHECK-SM80-NOF16-NEXT: min.f32 %r3, %r2, %r1;1632; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1633; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1634; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1635; CHECK-SM80-NOF16-NEXT: min.f32 %r6, %r5, %r4;1636; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1637; CHECK-SM80-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1638; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1639; CHECK-SM80-NOF16-NEXT: ret;1640 %x = call <2 x half> @llvm.minimumnum.v2f16(<2 x half> %a, <2 x half> %b)1641 ret <2 x half> %x1642}1643 1644; ---- maximumnum ----1645 1646define half @maximumnum_half(half %a, half %b) {1647; CHECK-NOF16-LABEL: maximumnum_half(1648; CHECK-NOF16: {1649; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;1650; CHECK-NOF16-NEXT: .reg .b32 %r<4>;1651; CHECK-NOF16-EMPTY:1652; CHECK-NOF16-NEXT: // %bb.0:1653; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];1654; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];1655; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1656; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1657; CHECK-NOF16-NEXT: max.f32 %r3, %r2, %r1;1658; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1659; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;1660; CHECK-NOF16-NEXT: ret;1661;1662; CHECK-F16-LABEL: maximumnum_half(1663; CHECK-F16: {1664; CHECK-F16-NEXT: .reg .b16 %rs<4>;1665; CHECK-F16-EMPTY:1666; CHECK-F16-NEXT: // %bb.0:1667; CHECK-F16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];1668; CHECK-F16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];1669; CHECK-F16-NEXT: max.f16 %rs3, %rs1, %rs2;1670; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;1671; CHECK-F16-NEXT: ret;1672;1673; CHECK-SM80-NOF16-LABEL: maximumnum_half(1674; CHECK-SM80-NOF16: {1675; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;1676; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;1677; CHECK-SM80-NOF16-EMPTY:1678; CHECK-SM80-NOF16-NEXT: // %bb.0:1679; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];1680; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];1681; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;1682; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;1683; CHECK-SM80-NOF16-NEXT: max.f32 %r3, %r2, %r1;1684; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1685; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;1686; CHECK-SM80-NOF16-NEXT: ret;1687 %x = call half @llvm.maximumnum.f16(half %a, half %b)1688 ret half %x1689}1690 1691define float @maximumnum_float(float %a, float %b) {1692; CHECK-LABEL: maximumnum_float(1693; CHECK: {1694; CHECK-NEXT: .reg .b32 %r<4>;1695; CHECK-EMPTY:1696; CHECK-NEXT: // %bb.0:1697; CHECK-NEXT: ld.param.b32 %r1, [maximumnum_float_param_0];1698; CHECK-NEXT: ld.param.b32 %r2, [maximumnum_float_param_1];1699; CHECK-NEXT: max.f32 %r3, %r1, %r2;1700; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1701; CHECK-NEXT: ret;1702 %x = call float @llvm.maximumnum.f32(float %a, float %b)1703 ret float %x1704}1705 1706define float @maximumnum_float_ftz(float %a, float %b) #1 {1707; CHECK-LABEL: maximumnum_float_ftz(1708; CHECK: {1709; CHECK-NEXT: .reg .b32 %r<4>;1710; CHECK-EMPTY:1711; CHECK-NEXT: // %bb.0:1712; CHECK-NEXT: ld.param.b32 %r1, [maximumnum_float_ftz_param_0];1713; CHECK-NEXT: ld.param.b32 %r2, [maximumnum_float_ftz_param_1];1714; CHECK-NEXT: max.ftz.f32 %r3, %r1, %r2;1715; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1716; CHECK-NEXT: ret;1717 %x = call float @llvm.maximumnum.f32(float %a, float %b)1718 ret float %x1719}1720 1721define double @maximumnum_double(double %a, double %b) {1722; CHECK-LABEL: maximumnum_double(1723; CHECK: {1724; CHECK-NEXT: .reg .b64 %rd<4>;1725; CHECK-EMPTY:1726; CHECK-NEXT: // %bb.0:1727; CHECK-NEXT: ld.param.b64 %rd1, [maximumnum_double_param_0];1728; CHECK-NEXT: ld.param.b64 %rd2, [maximumnum_double_param_1];1729; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;1730; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;1731; CHECK-NEXT: ret;1732 %x = call double @llvm.maximumnum.f64(double %a, double %b)1733 ret double %x1734}1735 1736define <2 x half> @maximumnum_v2half(<2 x half> %a, <2 x half> %b) {1737; CHECK-NOF16-LABEL: maximumnum_v2half(1738; CHECK-NOF16: {1739; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1740; CHECK-NOF16-NEXT: .reg .b32 %r<8>;1741; CHECK-NOF16-EMPTY:1742; CHECK-NOF16-NEXT: // %bb.0:1743; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximumnum_v2half_param_0];1744; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximumnum_v2half_param_1];1745; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1746; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1747; CHECK-NOF16-NEXT: max.f32 %r3, %r2, %r1;1748; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1749; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1750; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1751; CHECK-NOF16-NEXT: max.f32 %r6, %r5, %r4;1752; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1753; CHECK-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1754; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1755; CHECK-NOF16-NEXT: ret;1756;1757; CHECK-F16-LABEL: maximumnum_v2half(1758; CHECK-F16: {1759; CHECK-F16-NEXT: .reg .b32 %r<4>;1760; CHECK-F16-EMPTY:1761; CHECK-F16-NEXT: // %bb.0:1762; CHECK-F16-NEXT: ld.param.b32 %r1, [maximumnum_v2half_param_0];1763; CHECK-F16-NEXT: ld.param.b32 %r2, [maximumnum_v2half_param_1];1764; CHECK-F16-NEXT: max.f16x2 %r3, %r1, %r2;1765; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;1766; CHECK-F16-NEXT: ret;1767;1768; CHECK-SM80-NOF16-LABEL: maximumnum_v2half(1769; CHECK-SM80-NOF16: {1770; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;1771; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<8>;1772; CHECK-SM80-NOF16-EMPTY:1773; CHECK-SM80-NOF16-NEXT: // %bb.0:1774; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximumnum_v2half_param_0];1775; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximumnum_v2half_param_1];1776; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs4;1777; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1778; CHECK-SM80-NOF16-NEXT: max.f32 %r3, %r2, %r1;1779; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r3;1780; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;1781; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1782; CHECK-SM80-NOF16-NEXT: max.f32 %r6, %r5, %r4;1783; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r6;1784; CHECK-SM80-NOF16-NEXT: mov.b32 %r7, {%rs6, %rs5};1785; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0], %r7;1786; CHECK-SM80-NOF16-NEXT: ret;1787 %x = call <2 x half> @llvm.maximumnum.v2f16(<2 x half> %a, <2 x half> %b)1788 ret <2 x half> %x1789}1790 1791; ---- fma ----1792 1793define float @fma_float(float %a, float %b, float %c) {1794; CHECK-LABEL: fma_float(1795; CHECK: {1796; CHECK-NEXT: .reg .b32 %r<5>;1797; CHECK-EMPTY:1798; CHECK-NEXT: // %bb.0:1799; CHECK-NEXT: ld.param.b32 %r1, [fma_float_param_0];1800; CHECK-NEXT: ld.param.b32 %r2, [fma_float_param_1];1801; CHECK-NEXT: ld.param.b32 %r3, [fma_float_param_2];1802; CHECK-NEXT: fma.rn.f32 %r4, %r1, %r2, %r3;1803; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1804; CHECK-NEXT: ret;1805 %x = call float @llvm.fma.f32(float %a, float %b, float %c)1806 ret float %x1807}1808 1809define float @fma_float_ftz(float %a, float %b, float %c) #1 {1810; CHECK-LABEL: fma_float_ftz(1811; CHECK: {1812; CHECK-NEXT: .reg .b32 %r<5>;1813; CHECK-EMPTY:1814; CHECK-NEXT: // %bb.0:1815; CHECK-NEXT: ld.param.b32 %r1, [fma_float_ftz_param_0];1816; CHECK-NEXT: ld.param.b32 %r2, [fma_float_ftz_param_1];1817; CHECK-NEXT: ld.param.b32 %r3, [fma_float_ftz_param_2];1818; CHECK-NEXT: fma.rn.ftz.f32 %r4, %r1, %r2, %r3;1819; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1820; CHECK-NEXT: ret;1821 %x = call float @llvm.fma.f32(float %a, float %b, float %c)1822 ret float %x1823}1824 1825define double @fma_double(double %a, double %b, double %c) {1826; CHECK-LABEL: fma_double(1827; CHECK: {1828; CHECK-NEXT: .reg .b64 %rd<5>;1829; CHECK-EMPTY:1830; CHECK-NEXT: // %bb.0:1831; CHECK-NEXT: ld.param.b64 %rd1, [fma_double_param_0];1832; CHECK-NEXT: ld.param.b64 %rd2, [fma_double_param_1];1833; CHECK-NEXT: ld.param.b64 %rd3, [fma_double_param_2];1834; CHECK-NEXT: fma.rn.f64 %rd4, %rd1, %rd2, %rd3;1835; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;1836; CHECK-NEXT: ret;1837 %x = call double @llvm.fma.f64(double %a, double %b, double %c)1838 ret double %x1839}1840 1841attributes #0 = { nounwind readnone }1842attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }1843