1490 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -O1 | FileCheck %s --check-prefixes=CHECK,O1,O1-NO-MAD3; RUN: llc < %s -O1 -nvptx-mad-wide-opt | FileCheck %s --check-prefixes=CHECK,O1,O1-MAD4; RUN: llc < %s -O0 | FileCheck %s --check-prefixes=CHECK,O05 6target triple = "nvptx64-nvidia-cuda"7 8define i64 @t1(i32 %a, i32 %b, i64 %c) {9; O1-NO-MAD-LABEL: t1(10; O1-NO-MAD: {11; O1-NO-MAD-NEXT: .reg .b32 %r<3>;12; O1-NO-MAD-NEXT: .reg .b64 %rd<4>;13; O1-NO-MAD-EMPTY:14; O1-NO-MAD-NEXT: // %bb.0:15; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t1_param_0];16; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t1_param_1];17; O1-NO-MAD-NEXT: mul.wide.s32 %rd1, %r1, %r2;18; O1-NO-MAD-NEXT: ld.param.b64 %rd2, [t1_param_2];19; O1-NO-MAD-NEXT: add.s64 %rd3, %rd2, %rd1;20; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd3;21; O1-NO-MAD-NEXT: ret;22;23; O1-MAD-LABEL: t1(24; O1-MAD: {25; O1-MAD-NEXT: .reg .b32 %r<3>;26; O1-MAD-NEXT: .reg .b64 %rd<3>;27; O1-MAD-EMPTY:28; O1-MAD-NEXT: // %bb.0:29; O1-MAD-NEXT: ld.param.b32 %r1, [t1_param_0];30; O1-MAD-NEXT: ld.param.b32 %r2, [t1_param_1];31; O1-MAD-NEXT: ld.param.b64 %rd1, [t1_param_2];32; O1-MAD-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;33; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd2;34; O1-MAD-NEXT: ret;35;36; O0-LABEL: t1(37; O0: {38; O0-NEXT: .reg .b32 %r<4>;39; O0-NEXT: .reg .b64 %rd<4>;40; O0-EMPTY:41; O0-NEXT: // %bb.0:42; O0-NEXT: ld.param.b64 %rd1, [t1_param_2];43; O0-NEXT: ld.param.b32 %r2, [t1_param_1];44; O0-NEXT: ld.param.b32 %r1, [t1_param_0];45; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;46; O0-NEXT: cvt.s64.s32 %rd2, %r3;47; O0-NEXT: add.s64 %rd3, %rd1, %rd2;48; O0-NEXT: st.param.b64 [func_retval0], %rd3;49; O0-NEXT: ret;50 %mul = mul nsw i32 %a, %b51 %sext = sext i32 %mul to i6452 %add = add i64 %c, %sext53 ret i64 %add54}55 56define i64 @t2(i32 %a, i32 %b, i64 %c) {57; O1-NO-MAD-LABEL: t2(58; O1-NO-MAD: {59; O1-NO-MAD-NEXT: .reg .b32 %r<3>;60; O1-NO-MAD-NEXT: .reg .b64 %rd<4>;61; O1-NO-MAD-EMPTY:62; O1-NO-MAD-NEXT: // %bb.0:63; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t2_param_0];64; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t2_param_1];65; O1-NO-MAD-NEXT: mul.wide.s32 %rd1, %r1, %r2;66; O1-NO-MAD-NEXT: ld.param.b64 %rd2, [t2_param_2];67; O1-NO-MAD-NEXT: add.s64 %rd3, %rd1, %rd2;68; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd3;69; O1-NO-MAD-NEXT: ret;70;71; O1-MAD-LABEL: t2(72; O1-MAD: {73; O1-MAD-NEXT: .reg .b32 %r<3>;74; O1-MAD-NEXT: .reg .b64 %rd<3>;75; O1-MAD-EMPTY:76; O1-MAD-NEXT: // %bb.0:77; O1-MAD-NEXT: ld.param.b32 %r1, [t2_param_0];78; O1-MAD-NEXT: ld.param.b32 %r2, [t2_param_1];79; O1-MAD-NEXT: ld.param.b64 %rd1, [t2_param_2];80; O1-MAD-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;81; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd2;82; O1-MAD-NEXT: ret;83;84; O0-LABEL: t2(85; O0: {86; O0-NEXT: .reg .b32 %r<4>;87; O0-NEXT: .reg .b64 %rd<4>;88; O0-EMPTY:89; O0-NEXT: // %bb.0:90; O0-NEXT: ld.param.b64 %rd1, [t2_param_2];91; O0-NEXT: ld.param.b32 %r2, [t2_param_1];92; O0-NEXT: ld.param.b32 %r1, [t2_param_0];93; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;94; O0-NEXT: cvt.s64.s32 %rd2, %r3;95; O0-NEXT: add.s64 %rd3, %rd2, %rd1;96; O0-NEXT: st.param.b64 [func_retval0], %rd3;97; O0-NEXT: ret;98 %mul = mul nsw i32 %a, %b99 %sext = sext i32 %mul to i64100 %add = add i64 %sext, %c101 ret i64 %add102}103 104define i64 @t3(i32 %a, i32 %b) {105; O1-NO-MAD-LABEL: t3(106; O1-NO-MAD: {107; O1-NO-MAD-NEXT: .reg .b32 %r<3>;108; O1-NO-MAD-NEXT: .reg .b64 %rd<3>;109; O1-NO-MAD-EMPTY:110; O1-NO-MAD-NEXT: // %bb.0:111; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t3_param_0];112; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t3_param_1];113; O1-NO-MAD-NEXT: mul.wide.s32 %rd1, %r1, %r2;114; O1-NO-MAD-NEXT: add.s64 %rd2, %rd1, 1;115; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd2;116; O1-NO-MAD-NEXT: ret;117;118; O1-MAD-LABEL: t3(119; O1-MAD: {120; O1-MAD-NEXT: .reg .b32 %r<3>;121; O1-MAD-NEXT: .reg .b64 %rd<2>;122; O1-MAD-EMPTY:123; O1-MAD-NEXT: // %bb.0:124; O1-MAD-NEXT: ld.param.b32 %r1, [t3_param_0];125; O1-MAD-NEXT: ld.param.b32 %r2, [t3_param_1];126; O1-MAD-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;127; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd1;128; O1-MAD-NEXT: ret;129;130; O0-LABEL: t3(131; O0: {132; O0-NEXT: .reg .b32 %r<4>;133; O0-NEXT: .reg .b64 %rd<3>;134; O0-EMPTY:135; O0-NEXT: // %bb.0:136; O0-NEXT: ld.param.b32 %r2, [t3_param_1];137; O0-NEXT: ld.param.b32 %r1, [t3_param_0];138; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;139; O0-NEXT: cvt.s64.s32 %rd1, %r3;140; O0-NEXT: add.s64 %rd2, %rd1, 1;141; O0-NEXT: st.param.b64 [func_retval0], %rd2;142; O0-NEXT: ret;143 %mul = mul nsw i32 %a, %b144 %sext = sext i32 %mul to i64145 %add = add i64 1, %sext146 ret i64 %add147}148 149define i64 @t4(i32 %a, i64 %c) {150; O1-NO-MAD-LABEL: t4(151; O1-NO-MAD: {152; O1-NO-MAD-NEXT: .reg .b32 %r<2>;153; O1-NO-MAD-NEXT: .reg .b64 %rd<4>;154; O1-NO-MAD-EMPTY:155; O1-NO-MAD-NEXT: // %bb.0:156; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t4_param_0];157; O1-NO-MAD-NEXT: ld.param.b64 %rd1, [t4_param_1];158; O1-NO-MAD-NEXT: mul.wide.s32 %rd2, %r1, 3;159; O1-NO-MAD-NEXT: add.s64 %rd3, %rd1, %rd2;160; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd3;161; O1-NO-MAD-NEXT: ret;162;163; O1-MAD-LABEL: t4(164; O1-MAD: {165; O1-MAD-NEXT: .reg .b32 %r<2>;166; O1-MAD-NEXT: .reg .b64 %rd<3>;167; O1-MAD-EMPTY:168; O1-MAD-NEXT: // %bb.0:169; O1-MAD-NEXT: ld.param.b32 %r1, [t4_param_0];170; O1-MAD-NEXT: ld.param.b64 %rd1, [t4_param_1];171; O1-MAD-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;172; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd2;173; O1-MAD-NEXT: ret;174;175; O0-LABEL: t4(176; O0: {177; O0-NEXT: .reg .b32 %r<3>;178; O0-NEXT: .reg .b64 %rd<4>;179; O0-EMPTY:180; O0-NEXT: // %bb.0:181; O0-NEXT: ld.param.b64 %rd1, [t4_param_1];182; O0-NEXT: ld.param.b32 %r1, [t4_param_0];183; O0-NEXT: mul.lo.s32 %r2, %r1, 3;184; O0-NEXT: cvt.s64.s32 %rd2, %r2;185; O0-NEXT: add.s64 %rd3, %rd1, %rd2;186; O0-NEXT: st.param.b64 [func_retval0], %rd3;187; O0-NEXT: ret;188 %mul = mul nsw i32 %a, 3189 %sext = sext i32 %mul to i64190 %add = add i64 %c, %sext191 ret i64 %add192}193 194define i64 @t4_1(i32 %a, i64 %c) {195; O1-NO-MAD-LABEL: t4_1(196; O1-NO-MAD: {197; O1-NO-MAD-NEXT: .reg .b32 %r<2>;198; O1-NO-MAD-NEXT: .reg .b64 %rd<3>;199; O1-NO-MAD-EMPTY:200; O1-NO-MAD-NEXT: // %bb.0:201; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t4_1_param_0];202; O1-NO-MAD-NEXT: mul.wide.s32 %rd1, %r1, 3;203; O1-NO-MAD-NEXT: add.s64 %rd2, %rd1, 5;204; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd2;205; O1-NO-MAD-NEXT: ret;206;207; O1-MAD-LABEL: t4_1(208; O1-MAD: {209; O1-MAD-NEXT: .reg .b32 %r<2>;210; O1-MAD-NEXT: .reg .b64 %rd<2>;211; O1-MAD-EMPTY:212; O1-MAD-NEXT: // %bb.0:213; O1-MAD-NEXT: ld.param.b32 %r1, [t4_1_param_0];214; O1-MAD-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;215; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd1;216; O1-MAD-NEXT: ret;217;218; O0-LABEL: t4_1(219; O0: {220; O0-NEXT: .reg .b32 %r<3>;221; O0-NEXT: .reg .b64 %rd<3>;222; O0-EMPTY:223; O0-NEXT: // %bb.0:224; O0-NEXT: ld.param.b32 %r1, [t4_1_param_0];225; O0-NEXT: mul.lo.s32 %r2, %r1, 3;226; O0-NEXT: cvt.s64.s32 %rd1, %r2;227; O0-NEXT: add.s64 %rd2, %rd1, 5;228; O0-NEXT: st.param.b64 [func_retval0], %rd2;229; O0-NEXT: ret;230 %mul = mul nsw i32 %a, 3231 %sext = sext i32 %mul to i64232 %add = add i64 5, %sext233 ret i64 %add234}235 236define i64 @t5(i32 %a, i32 %b, i64 %c) {237; O1-NO-MAD-LABEL: t5(238; O1-NO-MAD: {239; O1-NO-MAD-NEXT: .reg .b32 %r<3>;240; O1-NO-MAD-NEXT: .reg .b64 %rd<4>;241; O1-NO-MAD-EMPTY:242; O1-NO-MAD-NEXT: // %bb.0:243; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t5_param_0];244; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t5_param_1];245; O1-NO-MAD-NEXT: mul.wide.u32 %rd1, %r1, %r2;246; O1-NO-MAD-NEXT: ld.param.b64 %rd2, [t5_param_2];247; O1-NO-MAD-NEXT: add.s64 %rd3, %rd2, %rd1;248; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd3;249; O1-NO-MAD-NEXT: ret;250;251; O1-MAD-LABEL: t5(252; O1-MAD: {253; O1-MAD-NEXT: .reg .b32 %r<3>;254; O1-MAD-NEXT: .reg .b64 %rd<3>;255; O1-MAD-EMPTY:256; O1-MAD-NEXT: // %bb.0:257; O1-MAD-NEXT: ld.param.b32 %r1, [t5_param_0];258; O1-MAD-NEXT: ld.param.b32 %r2, [t5_param_1];259; O1-MAD-NEXT: ld.param.b64 %rd1, [t5_param_2];260; O1-MAD-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;261; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd2;262; O1-MAD-NEXT: ret;263;264; O0-LABEL: t5(265; O0: {266; O0-NEXT: .reg .b32 %r<4>;267; O0-NEXT: .reg .b64 %rd<4>;268; O0-EMPTY:269; O0-NEXT: // %bb.0:270; O0-NEXT: ld.param.b64 %rd1, [t5_param_2];271; O0-NEXT: ld.param.b32 %r2, [t5_param_1];272; O0-NEXT: ld.param.b32 %r1, [t5_param_0];273; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;274; O0-NEXT: cvt.u64.u32 %rd2, %r3;275; O0-NEXT: add.s64 %rd3, %rd1, %rd2;276; O0-NEXT: st.param.b64 [func_retval0], %rd3;277; O0-NEXT: ret;278 %mul = mul nuw i32 %a, %b279 %zext = zext i32 %mul to i64280 %add = add i64 %c, %zext281 ret i64 %add282}283 284define i64 @t6(i32 %a, i32 %b, i64 %c) {285; O1-NO-MAD-LABEL: t6(286; O1-NO-MAD: {287; O1-NO-MAD-NEXT: .reg .b32 %r<3>;288; O1-NO-MAD-NEXT: .reg .b64 %rd<4>;289; O1-NO-MAD-EMPTY:290; O1-NO-MAD-NEXT: // %bb.0:291; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t6_param_0];292; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t6_param_1];293; O1-NO-MAD-NEXT: mul.wide.u32 %rd1, %r1, %r2;294; O1-NO-MAD-NEXT: ld.param.b64 %rd2, [t6_param_2];295; O1-NO-MAD-NEXT: add.s64 %rd3, %rd1, %rd2;296; O1-NO-MAD-NEXT: st.param.b64 [func_retval0], %rd3;297; O1-NO-MAD-NEXT: ret;298;299; O1-MAD-LABEL: t6(300; O1-MAD: {301; O1-MAD-NEXT: .reg .b32 %r<3>;302; O1-MAD-NEXT: .reg .b64 %rd<3>;303; O1-MAD-EMPTY:304; O1-MAD-NEXT: // %bb.0:305; O1-MAD-NEXT: ld.param.b32 %r1, [t6_param_0];306; O1-MAD-NEXT: ld.param.b32 %r2, [t6_param_1];307; O1-MAD-NEXT: ld.param.b64 %rd1, [t6_param_2];308; O1-MAD-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;309; O1-MAD-NEXT: st.param.b64 [func_retval0], %rd2;310; O1-MAD-NEXT: ret;311;312; O0-LABEL: t6(313; O0: {314; O0-NEXT: .reg .b32 %r<4>;315; O0-NEXT: .reg .b64 %rd<4>;316; O0-EMPTY:317; O0-NEXT: // %bb.0:318; O0-NEXT: ld.param.b64 %rd1, [t6_param_2];319; O0-NEXT: ld.param.b32 %r2, [t6_param_1];320; O0-NEXT: ld.param.b32 %r1, [t6_param_0];321; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;322; O0-NEXT: cvt.u64.u32 %rd2, %r3;323; O0-NEXT: add.s64 %rd3, %rd2, %rd1;324; O0-NEXT: st.param.b64 [func_retval0], %rd3;325; O0-NEXT: ret;326 %mul = mul nuw i32 %a, %b327 %zext = zext i32 %mul to i64328 %add = add i64 %zext, %c329 ret i64 %add330}331 332define i32 @t7(i16 %a, i16 %b) {333; O1-LABEL: t7(334; O1: {335; O1-NEXT: .reg .b16 %rs<4>;336; O1-NEXT: .reg .b32 %r<2>;337; O1-EMPTY:338; O1-NEXT: // %bb.0:339; O1-NEXT: ld.param.b16 %rs1, [t7_param_0];340; O1-NEXT: ld.param.b16 %rs2, [t7_param_1];341; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;342; O1-NEXT: cvt.u32.u16 %r1, %rs3;343; O1-NEXT: st.param.b32 [func_retval0], %r1;344; O1-NEXT: ret;345;346; O0-LABEL: t7(347; O0: {348; O0-NEXT: .reg .b16 %rs<4>;349; O0-NEXT: .reg .b32 %r<2>;350; O0-EMPTY:351; O0-NEXT: // %bb.0:352; O0-NEXT: ld.param.b16 %rs2, [t7_param_1];353; O0-NEXT: ld.param.b16 %rs1, [t7_param_0];354; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;355; O0-NEXT: cvt.u32.u16 %r1, %rs3;356; O0-NEXT: st.param.b32 [func_retval0], %r1;357; O0-NEXT: ret;358 %mul = mul i16 %a, %b359 %zext = zext i16 %mul to i32360 ret i32 %zext361}362 363define i32 @t8(i16 %a, i16 %b) {364; O1-LABEL: t8(365; O1: {366; O1-NEXT: .reg .b16 %rs<4>;367; O1-NEXT: .reg .b32 %r<2>;368; O1-EMPTY:369; O1-NEXT: // %bb.0:370; O1-NEXT: ld.param.b16 %rs1, [t8_param_0];371; O1-NEXT: ld.param.b16 %rs2, [t8_param_1];372; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;373; O1-NEXT: cvt.s32.s16 %r1, %rs3;374; O1-NEXT: st.param.b32 [func_retval0], %r1;375; O1-NEXT: ret;376;377; O0-LABEL: t8(378; O0: {379; O0-NEXT: .reg .b16 %rs<4>;380; O0-NEXT: .reg .b32 %r<2>;381; O0-EMPTY:382; O0-NEXT: // %bb.0:383; O0-NEXT: ld.param.b16 %rs2, [t8_param_1];384; O0-NEXT: ld.param.b16 %rs1, [t8_param_0];385; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;386; O0-NEXT: cvt.s32.s16 %r1, %rs3;387; O0-NEXT: st.param.b32 [func_retval0], %r1;388; O0-NEXT: ret;389 %mul = mul i16 %a, %b390 %sext = sext i16 %mul to i32391 ret i32 %sext392}393 394define i64 @t9(i32 %a, i32 %b) {395; O1-LABEL: t9(396; O1: {397; O1-NEXT: .reg .b32 %r<4>;398; O1-NEXT: .reg .b64 %rd<2>;399; O1-EMPTY:400; O1-NEXT: // %bb.0:401; O1-NEXT: ld.param.b32 %r1, [t9_param_0];402; O1-NEXT: ld.param.b32 %r2, [t9_param_1];403; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;404; O1-NEXT: cvt.u64.u32 %rd1, %r3;405; O1-NEXT: st.param.b64 [func_retval0], %rd1;406; O1-NEXT: ret;407;408; O0-LABEL: t9(409; O0: {410; O0-NEXT: .reg .b32 %r<4>;411; O0-NEXT: .reg .b64 %rd<2>;412; O0-EMPTY:413; O0-NEXT: // %bb.0:414; O0-NEXT: ld.param.b32 %r2, [t9_param_1];415; O0-NEXT: ld.param.b32 %r1, [t9_param_0];416; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;417; O0-NEXT: cvt.u64.u32 %rd1, %r3;418; O0-NEXT: st.param.b64 [func_retval0], %rd1;419; O0-NEXT: ret;420 %mul = mul i32 %a, %b421 %zext = zext i32 %mul to i64422 ret i64 %zext423}424 425define i64 @t10(i32 %a, i32 %b) {426; O1-LABEL: t10(427; O1: {428; O1-NEXT: .reg .b32 %r<4>;429; O1-NEXT: .reg .b64 %rd<2>;430; O1-EMPTY:431; O1-NEXT: // %bb.0:432; O1-NEXT: ld.param.b32 %r1, [t10_param_0];433; O1-NEXT: ld.param.b32 %r2, [t10_param_1];434; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;435; O1-NEXT: cvt.s64.s32 %rd1, %r3;436; O1-NEXT: st.param.b64 [func_retval0], %rd1;437; O1-NEXT: ret;438;439; O0-LABEL: t10(440; O0: {441; O0-NEXT: .reg .b32 %r<4>;442; O0-NEXT: .reg .b64 %rd<2>;443; O0-EMPTY:444; O0-NEXT: // %bb.0:445; O0-NEXT: ld.param.b32 %r2, [t10_param_1];446; O0-NEXT: ld.param.b32 %r1, [t10_param_0];447; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;448; O0-NEXT: cvt.s64.s32 %rd1, %r3;449; O0-NEXT: st.param.b64 [func_retval0], %rd1;450; O0-NEXT: ret;451 %mul = mul i32 %a, %b452 %sext = sext i32 %mul to i64453 ret i64 %sext454}455 456define i32 @t11(i16 %a, i16 %b) {457; O1-LABEL: t11(458; O1: {459; O1-NEXT: .reg .b16 %rs<4>;460; O1-NEXT: .reg .b32 %r<2>;461; O1-EMPTY:462; O1-NEXT: // %bb.0:463; O1-NEXT: ld.param.b16 %rs1, [t11_param_0];464; O1-NEXT: ld.param.b16 %rs2, [t11_param_1];465; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;466; O1-NEXT: cvt.u32.u16 %r1, %rs3;467; O1-NEXT: st.param.b32 [func_retval0], %r1;468; O1-NEXT: ret;469;470; O0-LABEL: t11(471; O0: {472; O0-NEXT: .reg .b16 %rs<4>;473; O0-NEXT: .reg .b32 %r<2>;474; O0-EMPTY:475; O0-NEXT: // %bb.0:476; O0-NEXT: ld.param.b16 %rs2, [t11_param_1];477; O0-NEXT: ld.param.b16 %rs1, [t11_param_0];478; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;479; O0-NEXT: cvt.u32.u16 %r1, %rs3;480; O0-NEXT: st.param.b32 [func_retval0], %r1;481; O0-NEXT: ret;482 %mul = mul nsw i16 %a, %b483 %zext = zext i16 %mul to i32484 ret i32 %zext485}486 487define i32 @t12(i16 %a, i16 %b) {488; O1-LABEL: t12(489; O1: {490; O1-NEXT: .reg .b16 %rs<3>;491; O1-NEXT: .reg .b32 %r<2>;492; O1-EMPTY:493; O1-NEXT: // %bb.0:494; O1-NEXT: ld.param.b16 %rs1, [t12_param_0];495; O1-NEXT: ld.param.b16 %rs2, [t12_param_1];496; O1-NEXT: mul.wide.s16 %r1, %rs1, %rs2;497; O1-NEXT: st.param.b32 [func_retval0], %r1;498; O1-NEXT: ret;499;500; O0-LABEL: t12(501; O0: {502; O0-NEXT: .reg .b16 %rs<4>;503; O0-NEXT: .reg .b32 %r<2>;504; O0-EMPTY:505; O0-NEXT: // %bb.0:506; O0-NEXT: ld.param.b16 %rs2, [t12_param_1];507; O0-NEXT: ld.param.b16 %rs1, [t12_param_0];508; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;509; O0-NEXT: cvt.s32.s16 %r1, %rs3;510; O0-NEXT: st.param.b32 [func_retval0], %r1;511; O0-NEXT: ret;512 %mul = mul nsw i16 %a, %b513 %sext = sext i16 %mul to i32514 ret i32 %sext515}516 517define i64 @t13(i32 %a, i32 %b) {518; O1-LABEL: t13(519; O1: {520; O1-NEXT: .reg .b32 %r<4>;521; O1-NEXT: .reg .b64 %rd<2>;522; O1-EMPTY:523; O1-NEXT: // %bb.0:524; O1-NEXT: ld.param.b32 %r1, [t13_param_0];525; O1-NEXT: ld.param.b32 %r2, [t13_param_1];526; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;527; O1-NEXT: cvt.u64.u32 %rd1, %r3;528; O1-NEXT: st.param.b64 [func_retval0], %rd1;529; O1-NEXT: ret;530;531; O0-LABEL: t13(532; O0: {533; O0-NEXT: .reg .b32 %r<4>;534; O0-NEXT: .reg .b64 %rd<2>;535; O0-EMPTY:536; O0-NEXT: // %bb.0:537; O0-NEXT: ld.param.b32 %r2, [t13_param_1];538; O0-NEXT: ld.param.b32 %r1, [t13_param_0];539; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;540; O0-NEXT: cvt.u64.u32 %rd1, %r3;541; O0-NEXT: st.param.b64 [func_retval0], %rd1;542; O0-NEXT: ret;543 %mul = mul nsw i32 %a, %b544 %zext = zext i32 %mul to i64545 ret i64 %zext546}547 548define i64 @t14(i32 %a, i32 %b) {549; O1-LABEL: t14(550; O1: {551; O1-NEXT: .reg .b32 %r<3>;552; O1-NEXT: .reg .b64 %rd<2>;553; O1-EMPTY:554; O1-NEXT: // %bb.0:555; O1-NEXT: ld.param.b32 %r1, [t14_param_0];556; O1-NEXT: ld.param.b32 %r2, [t14_param_1];557; O1-NEXT: mul.wide.s32 %rd1, %r1, %r2;558; O1-NEXT: st.param.b64 [func_retval0], %rd1;559; O1-NEXT: ret;560;561; O0-LABEL: t14(562; O0: {563; O0-NEXT: .reg .b32 %r<4>;564; O0-NEXT: .reg .b64 %rd<2>;565; O0-EMPTY:566; O0-NEXT: // %bb.0:567; O0-NEXT: ld.param.b32 %r2, [t14_param_1];568; O0-NEXT: ld.param.b32 %r1, [t14_param_0];569; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;570; O0-NEXT: cvt.s64.s32 %rd1, %r3;571; O0-NEXT: st.param.b64 [func_retval0], %rd1;572; O0-NEXT: ret;573 %mul = mul nsw i32 %a, %b574 %sext = sext i32 %mul to i64575 ret i64 %sext576}577 578define i32 @t15(i16 %a, i16 %b) {579; O1-LABEL: t15(580; O1: {581; O1-NEXT: .reg .b16 %rs<3>;582; O1-NEXT: .reg .b32 %r<2>;583; O1-EMPTY:584; O1-NEXT: // %bb.0:585; O1-NEXT: ld.param.b16 %rs1, [t15_param_0];586; O1-NEXT: ld.param.b16 %rs2, [t15_param_1];587; O1-NEXT: mul.wide.u16 %r1, %rs1, %rs2;588; O1-NEXT: st.param.b32 [func_retval0], %r1;589; O1-NEXT: ret;590;591; O0-LABEL: t15(592; O0: {593; O0-NEXT: .reg .b16 %rs<4>;594; O0-NEXT: .reg .b32 %r<2>;595; O0-EMPTY:596; O0-NEXT: // %bb.0:597; O0-NEXT: ld.param.b16 %rs2, [t15_param_1];598; O0-NEXT: ld.param.b16 %rs1, [t15_param_0];599; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;600; O0-NEXT: cvt.u32.u16 %r1, %rs3;601; O0-NEXT: st.param.b32 [func_retval0], %r1;602; O0-NEXT: ret;603 %mul = mul nuw i16 %a, %b604 %zext = zext i16 %mul to i32605 ret i32 %zext606}607 608define i32 @t16(i16 %a, i16 %b) {609; O1-LABEL: t16(610; O1: {611; O1-NEXT: .reg .b16 %rs<4>;612; O1-NEXT: .reg .b32 %r<2>;613; O1-EMPTY:614; O1-NEXT: // %bb.0:615; O1-NEXT: ld.param.b16 %rs1, [t16_param_0];616; O1-NEXT: ld.param.b16 %rs2, [t16_param_1];617; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;618; O1-NEXT: cvt.s32.s16 %r1, %rs3;619; O1-NEXT: st.param.b32 [func_retval0], %r1;620; O1-NEXT: ret;621;622; O0-LABEL: t16(623; O0: {624; O0-NEXT: .reg .b16 %rs<4>;625; O0-NEXT: .reg .b32 %r<2>;626; O0-EMPTY:627; O0-NEXT: // %bb.0:628; O0-NEXT: ld.param.b16 %rs2, [t16_param_1];629; O0-NEXT: ld.param.b16 %rs1, [t16_param_0];630; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;631; O0-NEXT: cvt.s32.s16 %r1, %rs3;632; O0-NEXT: st.param.b32 [func_retval0], %r1;633; O0-NEXT: ret;634 %mul = mul nuw i16 %a, %b635 %sext = sext i16 %mul to i32636 ret i32 %sext637}638 639define i64 @t17(i32 %a, i32 %b) {640; O1-LABEL: t17(641; O1: {642; O1-NEXT: .reg .b32 %r<3>;643; O1-NEXT: .reg .b64 %rd<2>;644; O1-EMPTY:645; O1-NEXT: // %bb.0:646; O1-NEXT: ld.param.b32 %r1, [t17_param_0];647; O1-NEXT: ld.param.b32 %r2, [t17_param_1];648; O1-NEXT: mul.wide.u32 %rd1, %r1, %r2;649; O1-NEXT: st.param.b64 [func_retval0], %rd1;650; O1-NEXT: ret;651;652; O0-LABEL: t17(653; O0: {654; O0-NEXT: .reg .b32 %r<4>;655; O0-NEXT: .reg .b64 %rd<2>;656; O0-EMPTY:657; O0-NEXT: // %bb.0:658; O0-NEXT: ld.param.b32 %r2, [t17_param_1];659; O0-NEXT: ld.param.b32 %r1, [t17_param_0];660; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;661; O0-NEXT: cvt.u64.u32 %rd1, %r3;662; O0-NEXT: st.param.b64 [func_retval0], %rd1;663; O0-NEXT: ret;664 %mul = mul nuw i32 %a, %b665 %zext = zext i32 %mul to i64666 ret i64 %zext667}668 669define i64 @t18(i32 %a, i32 %b) {670; O1-LABEL: t18(671; O1: {672; O1-NEXT: .reg .b32 %r<4>;673; O1-NEXT: .reg .b64 %rd<2>;674; O1-EMPTY:675; O1-NEXT: // %bb.0:676; O1-NEXT: ld.param.b32 %r1, [t18_param_0];677; O1-NEXT: ld.param.b32 %r2, [t18_param_1];678; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;679; O1-NEXT: cvt.s64.s32 %rd1, %r3;680; O1-NEXT: st.param.b64 [func_retval0], %rd1;681; O1-NEXT: ret;682;683; O0-LABEL: t18(684; O0: {685; O0-NEXT: .reg .b32 %r<4>;686; O0-NEXT: .reg .b64 %rd<2>;687; O0-EMPTY:688; O0-NEXT: // %bb.0:689; O0-NEXT: ld.param.b32 %r2, [t18_param_1];690; O0-NEXT: ld.param.b32 %r1, [t18_param_0];691; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;692; O0-NEXT: cvt.s64.s32 %rd1, %r3;693; O0-NEXT: st.param.b64 [func_retval0], %rd1;694; O0-NEXT: ret;695 %mul = mul nuw i32 %a, %b696 %sext = sext i32 %mul to i64697 ret i64 %sext698}699 700define i32 @t19(i16 %a, i16 %b) {701; O1-LABEL: t19(702; O1: {703; O1-NEXT: .reg .b16 %rs<4>;704; O1-NEXT: .reg .b32 %r<2>;705; O1-EMPTY:706; O1-NEXT: // %bb.0:707; O1-NEXT: ld.param.b16 %rs1, [t19_param_0];708; O1-NEXT: ld.param.b16 %rs2, [t19_param_1];709; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;710; O1-NEXT: cvt.u32.u16 %r1, %rs3;711; O1-NEXT: st.param.b32 [func_retval0], %r1;712; O1-NEXT: ret;713;714; O0-LABEL: t19(715; O0: {716; O0-NEXT: .reg .b16 %rs<4>;717; O0-NEXT: .reg .b32 %r<2>;718; O0-EMPTY:719; O0-NEXT: // %bb.0:720; O0-NEXT: ld.param.b16 %rs2, [t19_param_1];721; O0-NEXT: ld.param.b16 %rs1, [t19_param_0];722; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;723; O0-NEXT: cvt.u32.u16 %r1, %rs3;724; O0-NEXT: st.param.b32 [func_retval0], %r1;725; O0-NEXT: ret;726 %mul = mul i16 %a, %b727 %zext = zext i16 %mul to i32728 ret i32 %zext729}730 731define i32 @t20(i16 %a) {732; CHECK-LABEL: t20(733; CHECK: {734; CHECK-NEXT: .reg .b16 %rs<3>;735; CHECK-NEXT: .reg .b32 %r<2>;736; CHECK-EMPTY:737; CHECK-NEXT: // %bb.0:738; CHECK-NEXT: ld.param.b16 %rs1, [t20_param_0];739; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;740; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;741; CHECK-NEXT: st.param.b32 [func_retval0], %r1;742; CHECK-NEXT: ret;743 %mul = shl i16 %a, 4744 %sext = sext i16 %mul to i32745 ret i32 %sext746}747 748define i64 @t21(i32 %a) {749; CHECK-LABEL: t21(750; CHECK: {751; CHECK-NEXT: .reg .b32 %r<3>;752; CHECK-NEXT: .reg .b64 %rd<2>;753; CHECK-EMPTY:754; CHECK-NEXT: // %bb.0:755; CHECK-NEXT: ld.param.b32 %r1, [t21_param_0];756; CHECK-NEXT: shl.b32 %r2, %r1, 4;757; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;758; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;759; CHECK-NEXT: ret;760 %mul = shl i32 %a, 4761 %zext = zext i32 %mul to i64762 ret i64 %zext763}764 765define i64 @t22(i32 %a) {766; CHECK-LABEL: t22(767; CHECK: {768; CHECK-NEXT: .reg .b32 %r<3>;769; CHECK-NEXT: .reg .b64 %rd<2>;770; CHECK-EMPTY:771; CHECK-NEXT: // %bb.0:772; CHECK-NEXT: ld.param.b32 %r1, [t22_param_0];773; CHECK-NEXT: shl.b32 %r2, %r1, 4;774; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;775; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;776; CHECK-NEXT: ret;777 %mul = shl i32 %a, 4778 %sext = sext i32 %mul to i64779 ret i64 %sext780}781 782define i32 @t23(i16 %a, i16 %b) {783; CHECK-LABEL: t23(784; CHECK: {785; CHECK-NEXT: .reg .b16 %rs<3>;786; CHECK-NEXT: .reg .b32 %r<2>;787; CHECK-EMPTY:788; CHECK-NEXT: // %bb.0:789; CHECK-NEXT: ld.param.b16 %rs1, [t23_param_0];790; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;791; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;792; CHECK-NEXT: st.param.b32 [func_retval0], %r1;793; CHECK-NEXT: ret;794 %mul = shl nsw i16 %a, 4795 %zext = zext i16 %mul to i32796 ret i32 %zext797}798 799define i32 @t24(i16 %a, i16 %b) {800; O1-LABEL: t24(801; O1: {802; O1-NEXT: .reg .b16 %rs<2>;803; O1-NEXT: .reg .b32 %r<2>;804; O1-EMPTY:805; O1-NEXT: // %bb.0:806; O1-NEXT: ld.param.b16 %rs1, [t24_param_0];807; O1-NEXT: mul.wide.s16 %r1, %rs1, 16;808; O1-NEXT: st.param.b32 [func_retval0], %r1;809; O1-NEXT: ret;810;811; O0-LABEL: t24(812; O0: {813; O0-NEXT: .reg .b16 %rs<3>;814; O0-NEXT: .reg .b32 %r<2>;815; O0-EMPTY:816; O0-NEXT: // %bb.0:817; O0-NEXT: ld.param.b16 %rs1, [t24_param_0];818; O0-NEXT: shl.b16 %rs2, %rs1, 4;819; O0-NEXT: cvt.s32.s16 %r1, %rs2;820; O0-NEXT: st.param.b32 [func_retval0], %r1;821; O0-NEXT: ret;822 %mul = shl nsw i16 %a, 4823 %sext = sext i16 %mul to i32824 ret i32 %sext825}826 827define i64 @t25(i32 %a) {828; CHECK-LABEL: t25(829; CHECK: {830; CHECK-NEXT: .reg .b32 %r<3>;831; CHECK-NEXT: .reg .b64 %rd<2>;832; CHECK-EMPTY:833; CHECK-NEXT: // %bb.0:834; CHECK-NEXT: ld.param.b32 %r1, [t25_param_0];835; CHECK-NEXT: shl.b32 %r2, %r1, 4;836; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;837; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;838; CHECK-NEXT: ret;839 %mul = shl nsw i32 %a, 4840 %zext = zext i32 %mul to i64841 ret i64 %zext842}843 844define i64 @t26(i32 %a) {845; O1-LABEL: t26(846; O1: {847; O1-NEXT: .reg .b32 %r<2>;848; O1-NEXT: .reg .b64 %rd<2>;849; O1-EMPTY:850; O1-NEXT: // %bb.0:851; O1-NEXT: ld.param.b32 %r1, [t26_param_0];852; O1-NEXT: mul.wide.s32 %rd1, %r1, 16;853; O1-NEXT: st.param.b64 [func_retval0], %rd1;854; O1-NEXT: ret;855;856; O0-LABEL: t26(857; O0: {858; O0-NEXT: .reg .b32 %r<3>;859; O0-NEXT: .reg .b64 %rd<2>;860; O0-EMPTY:861; O0-NEXT: // %bb.0:862; O0-NEXT: ld.param.b32 %r1, [t26_param_0];863; O0-NEXT: shl.b32 %r2, %r1, 4;864; O0-NEXT: cvt.s64.s32 %rd1, %r2;865; O0-NEXT: st.param.b64 [func_retval0], %rd1;866; O0-NEXT: ret;867 %mul = shl nsw i32 %a, 4868 %sext = sext i32 %mul to i64869 ret i64 %sext870}871 872define i32 @t27(i16 %a, i16 %b) {873; O1-LABEL: t27(874; O1: {875; O1-NEXT: .reg .b16 %rs<2>;876; O1-NEXT: .reg .b32 %r<2>;877; O1-EMPTY:878; O1-NEXT: // %bb.0:879; O1-NEXT: ld.param.b16 %rs1, [t27_param_0];880; O1-NEXT: mul.wide.u16 %r1, %rs1, 16;881; O1-NEXT: st.param.b32 [func_retval0], %r1;882; O1-NEXT: ret;883;884; O0-LABEL: t27(885; O0: {886; O0-NEXT: .reg .b16 %rs<3>;887; O0-NEXT: .reg .b32 %r<2>;888; O0-EMPTY:889; O0-NEXT: // %bb.0:890; O0-NEXT: ld.param.b16 %rs1, [t27_param_0];891; O0-NEXT: shl.b16 %rs2, %rs1, 4;892; O0-NEXT: cvt.u32.u16 %r1, %rs2;893; O0-NEXT: st.param.b32 [func_retval0], %r1;894; O0-NEXT: ret;895 %mul = shl nuw i16 %a, 4896 %zext = zext i16 %mul to i32897 ret i32 %zext898}899 900define i32 @t28(i16 %a, i16 %b) {901; CHECK-LABEL: t28(902; CHECK: {903; CHECK-NEXT: .reg .b16 %rs<3>;904; CHECK-NEXT: .reg .b32 %r<2>;905; CHECK-EMPTY:906; CHECK-NEXT: // %bb.0:907; CHECK-NEXT: ld.param.b16 %rs1, [t28_param_0];908; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;909; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;910; CHECK-NEXT: st.param.b32 [func_retval0], %r1;911; CHECK-NEXT: ret;912 %mul = shl nuw i16 %a, 4913 %sext = sext i16 %mul to i32914 ret i32 %sext915}916 917define i64 @t29(i32 %a) {918; O1-LABEL: t29(919; O1: {920; O1-NEXT: .reg .b32 %r<2>;921; O1-NEXT: .reg .b64 %rd<2>;922; O1-EMPTY:923; O1-NEXT: // %bb.0:924; O1-NEXT: ld.param.b32 %r1, [t29_param_0];925; O1-NEXT: mul.wide.u32 %rd1, %r1, 16;926; O1-NEXT: st.param.b64 [func_retval0], %rd1;927; O1-NEXT: ret;928;929; O0-LABEL: t29(930; O0: {931; O0-NEXT: .reg .b32 %r<3>;932; O0-NEXT: .reg .b64 %rd<2>;933; O0-EMPTY:934; O0-NEXT: // %bb.0:935; O0-NEXT: ld.param.b32 %r1, [t29_param_0];936; O0-NEXT: shl.b32 %r2, %r1, 4;937; O0-NEXT: cvt.u64.u32 %rd1, %r2;938; O0-NEXT: st.param.b64 [func_retval0], %rd1;939; O0-NEXT: ret;940 %mul = shl nuw i32 %a, 4941 %zext = zext i32 %mul to i64942 ret i64 %zext943}944 945define i64 @t30(i32 %a) {946; CHECK-LABEL: t30(947; CHECK: {948; CHECK-NEXT: .reg .b32 %r<3>;949; CHECK-NEXT: .reg .b64 %rd<2>;950; CHECK-EMPTY:951; CHECK-NEXT: // %bb.0:952; CHECK-NEXT: ld.param.b32 %r1, [t30_param_0];953; CHECK-NEXT: shl.b32 %r2, %r1, 4;954; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;955; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;956; CHECK-NEXT: ret;957 %mul = shl nuw i32 %a, 4958 %sext = sext i32 %mul to i64959 ret i64 %sext960}961 962define i64 @t31(i32 %a, i32 %b) {963; O1-LABEL: t31(964; O1: {965; O1-NEXT: .reg .b32 %r<4>;966; O1-NEXT: .reg .b64 %rd<2>;967; O1-EMPTY:968; O1-NEXT: // %bb.0:969; O1-NEXT: ld.param.b32 %r1, [t31_param_0];970; O1-NEXT: ld.param.b32 %r2, [t31_param_1];971; O1-NEXT: shl.b32 %r3, %r1, %r2;972; O1-NEXT: cvt.s64.s32 %rd1, %r3;973; O1-NEXT: st.param.b64 [func_retval0], %rd1;974; O1-NEXT: ret;975;976; O0-LABEL: t31(977; O0: {978; O0-NEXT: .reg .b32 %r<4>;979; O0-NEXT: .reg .b64 %rd<2>;980; O0-EMPTY:981; O0-NEXT: // %bb.0:982; O0-NEXT: ld.param.b32 %r2, [t31_param_1];983; O0-NEXT: ld.param.b32 %r1, [t31_param_0];984; O0-NEXT: shl.b32 %r3, %r1, %r2;985; O0-NEXT: cvt.s64.s32 %rd1, %r3;986; O0-NEXT: st.param.b64 [func_retval0], %rd1;987; O0-NEXT: ret;988 %mul = shl nuw i32 %a, %b989 %sext = sext i32 %mul to i64990 ret i64 %sext991}992 993define i32 @t32(i16 %a, i16 %b, i32 %c) {994; O1-NO-MAD-LABEL: t32(995; O1-NO-MAD: {996; O1-NO-MAD-NEXT: .reg .b16 %rs<3>;997; O1-NO-MAD-NEXT: .reg .b32 %r<4>;998; O1-NO-MAD-EMPTY:999; O1-NO-MAD-NEXT: // %bb.0:1000; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t32_param_0];1001; O1-NO-MAD-NEXT: ld.param.b16 %rs2, [t32_param_1];1002; O1-NO-MAD-NEXT: mul.wide.s16 %r1, %rs1, %rs2;1003; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t32_param_2];1004; O1-NO-MAD-NEXT: add.s32 %r3, %r2, %r1;1005; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r3;1006; O1-NO-MAD-NEXT: ret;1007;1008; O1-MAD-LABEL: t32(1009; O1-MAD: {1010; O1-MAD-NEXT: .reg .b16 %rs<3>;1011; O1-MAD-NEXT: .reg .b32 %r<3>;1012; O1-MAD-EMPTY:1013; O1-MAD-NEXT: // %bb.0:1014; O1-MAD-NEXT: ld.param.b16 %rs1, [t32_param_0];1015; O1-MAD-NEXT: ld.param.b16 %rs2, [t32_param_1];1016; O1-MAD-NEXT: ld.param.b32 %r1, [t32_param_2];1017; O1-MAD-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;1018; O1-MAD-NEXT: st.param.b32 [func_retval0], %r2;1019; O1-MAD-NEXT: ret;1020;1021; O0-LABEL: t32(1022; O0: {1023; O0-NEXT: .reg .b16 %rs<4>;1024; O0-NEXT: .reg .b32 %r<4>;1025; O0-EMPTY:1026; O0-NEXT: // %bb.0:1027; O0-NEXT: ld.param.b32 %r1, [t32_param_2];1028; O0-NEXT: ld.param.b16 %rs2, [t32_param_1];1029; O0-NEXT: ld.param.b16 %rs1, [t32_param_0];1030; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1031; O0-NEXT: cvt.s32.s16 %r2, %rs3;1032; O0-NEXT: add.s32 %r3, %r1, %r2;1033; O0-NEXT: st.param.b32 [func_retval0], %r3;1034; O0-NEXT: ret;1035 %mul = mul nsw i16 %a, %b1036 %sext = sext i16 %mul to i321037 %add = add i32 %c, %sext1038 ret i32 %add1039}1040 1041define i32 @t33(i16 %a, i16 %b, i32 %c) {1042; O1-NO-MAD-LABEL: t33(1043; O1-NO-MAD: {1044; O1-NO-MAD-NEXT: .reg .b16 %rs<3>;1045; O1-NO-MAD-NEXT: .reg .b32 %r<4>;1046; O1-NO-MAD-EMPTY:1047; O1-NO-MAD-NEXT: // %bb.0:1048; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t33_param_0];1049; O1-NO-MAD-NEXT: ld.param.b16 %rs2, [t33_param_1];1050; O1-NO-MAD-NEXT: mul.wide.s16 %r1, %rs1, %rs2;1051; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t33_param_2];1052; O1-NO-MAD-NEXT: add.s32 %r3, %r2, %r1;1053; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r3;1054; O1-NO-MAD-NEXT: ret;1055;1056; O1-MAD-LABEL: t33(1057; O1-MAD: {1058; O1-MAD-NEXT: .reg .b16 %rs<3>;1059; O1-MAD-NEXT: .reg .b32 %r<3>;1060; O1-MAD-EMPTY:1061; O1-MAD-NEXT: // %bb.0:1062; O1-MAD-NEXT: ld.param.b16 %rs1, [t33_param_0];1063; O1-MAD-NEXT: ld.param.b16 %rs2, [t33_param_1];1064; O1-MAD-NEXT: ld.param.b32 %r1, [t33_param_2];1065; O1-MAD-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;1066; O1-MAD-NEXT: st.param.b32 [func_retval0], %r2;1067; O1-MAD-NEXT: ret;1068;1069; O0-LABEL: t33(1070; O0: {1071; O0-NEXT: .reg .b16 %rs<4>;1072; O0-NEXT: .reg .b32 %r<4>;1073; O0-EMPTY:1074; O0-NEXT: // %bb.0:1075; O0-NEXT: ld.param.b32 %r1, [t33_param_2];1076; O0-NEXT: ld.param.b16 %rs2, [t33_param_1];1077; O0-NEXT: ld.param.b16 %rs1, [t33_param_0];1078; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1079; O0-NEXT: cvt.s32.s16 %r2, %rs3;1080; O0-NEXT: add.s32 %r3, %r1, %r2;1081; O0-NEXT: st.param.b32 [func_retval0], %r3;1082; O0-NEXT: ret;1083 %mul = mul nsw i16 %a, %b1084 %sext = sext i16 %mul to i321085 %add = add i32 %c, %sext1086 ret i32 %add1087}1088 1089define i32 @t34(i16 %a, i16 %b) {1090; O1-NO-MAD-LABEL: t34(1091; O1-NO-MAD: {1092; O1-NO-MAD-NEXT: .reg .b16 %rs<3>;1093; O1-NO-MAD-NEXT: .reg .b32 %r<3>;1094; O1-NO-MAD-EMPTY:1095; O1-NO-MAD-NEXT: // %bb.0:1096; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t34_param_0];1097; O1-NO-MAD-NEXT: ld.param.b16 %rs2, [t34_param_1];1098; O1-NO-MAD-NEXT: mul.wide.s16 %r1, %rs1, %rs2;1099; O1-NO-MAD-NEXT: add.s32 %r2, %r1, 1;1100; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r2;1101; O1-NO-MAD-NEXT: ret;1102;1103; O1-MAD-LABEL: t34(1104; O1-MAD: {1105; O1-MAD-NEXT: .reg .b16 %rs<3>;1106; O1-MAD-NEXT: .reg .b32 %r<2>;1107; O1-MAD-EMPTY:1108; O1-MAD-NEXT: // %bb.0:1109; O1-MAD-NEXT: ld.param.b16 %rs1, [t34_param_0];1110; O1-MAD-NEXT: ld.param.b16 %rs2, [t34_param_1];1111; O1-MAD-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;1112; O1-MAD-NEXT: st.param.b32 [func_retval0], %r1;1113; O1-MAD-NEXT: ret;1114;1115; O0-LABEL: t34(1116; O0: {1117; O0-NEXT: .reg .b16 %rs<4>;1118; O0-NEXT: .reg .b32 %r<3>;1119; O0-EMPTY:1120; O0-NEXT: // %bb.0:1121; O0-NEXT: ld.param.b16 %rs2, [t34_param_1];1122; O0-NEXT: ld.param.b16 %rs1, [t34_param_0];1123; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1124; O0-NEXT: cvt.s32.s16 %r1, %rs3;1125; O0-NEXT: add.s32 %r2, %r1, 1;1126; O0-NEXT: st.param.b32 [func_retval0], %r2;1127; O0-NEXT: ret;1128 %mul = mul nsw i16 %a, %b1129 %sext = sext i16 %mul to i321130 %add = add i32 1, %sext1131 ret i32 %add1132}1133 1134define i32 @t35(i16 %a, i32 %c) {1135; O1-NO-MAD-LABEL: t35(1136; O1-NO-MAD: {1137; O1-NO-MAD-NEXT: .reg .b16 %rs<2>;1138; O1-NO-MAD-NEXT: .reg .b32 %r<4>;1139; O1-NO-MAD-EMPTY:1140; O1-NO-MAD-NEXT: // %bb.0:1141; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t35_param_0];1142; O1-NO-MAD-NEXT: ld.param.b32 %r1, [t35_param_1];1143; O1-NO-MAD-NEXT: mul.wide.s16 %r2, %rs1, 3;1144; O1-NO-MAD-NEXT: add.s32 %r3, %r1, %r2;1145; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r3;1146; O1-NO-MAD-NEXT: ret;1147;1148; O1-MAD-LABEL: t35(1149; O1-MAD: {1150; O1-MAD-NEXT: .reg .b16 %rs<2>;1151; O1-MAD-NEXT: .reg .b32 %r<3>;1152; O1-MAD-EMPTY:1153; O1-MAD-NEXT: // %bb.0:1154; O1-MAD-NEXT: ld.param.b16 %rs1, [t35_param_0];1155; O1-MAD-NEXT: ld.param.b32 %r1, [t35_param_1];1156; O1-MAD-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;1157; O1-MAD-NEXT: st.param.b32 [func_retval0], %r2;1158; O1-MAD-NEXT: ret;1159;1160; O0-LABEL: t35(1161; O0: {1162; O0-NEXT: .reg .b16 %rs<3>;1163; O0-NEXT: .reg .b32 %r<4>;1164; O0-EMPTY:1165; O0-NEXT: // %bb.0:1166; O0-NEXT: ld.param.b32 %r1, [t35_param_1];1167; O0-NEXT: ld.param.b16 %rs1, [t35_param_0];1168; O0-NEXT: mul.lo.s16 %rs2, %rs1, 3;1169; O0-NEXT: cvt.s32.s16 %r2, %rs2;1170; O0-NEXT: add.s32 %r3, %r1, %r2;1171; O0-NEXT: st.param.b32 [func_retval0], %r3;1172; O0-NEXT: ret;1173 %mul = mul nsw i16 %a, 31174 %sext = sext i16 %mul to i321175 %add = add i32 %c, %sext1176 ret i32 %add1177}1178 1179define i32 @t36(i16 %a, i32 %c) {1180; O1-NO-MAD-LABEL: t36(1181; O1-NO-MAD: {1182; O1-NO-MAD-NEXT: .reg .b16 %rs<2>;1183; O1-NO-MAD-NEXT: .reg .b32 %r<3>;1184; O1-NO-MAD-EMPTY:1185; O1-NO-MAD-NEXT: // %bb.0:1186; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t36_param_0];1187; O1-NO-MAD-NEXT: mul.wide.s16 %r1, %rs1, 3;1188; O1-NO-MAD-NEXT: add.s32 %r2, %r1, 5;1189; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r2;1190; O1-NO-MAD-NEXT: ret;1191;1192; O1-MAD-LABEL: t36(1193; O1-MAD: {1194; O1-MAD-NEXT: .reg .b16 %rs<2>;1195; O1-MAD-NEXT: .reg .b32 %r<2>;1196; O1-MAD-EMPTY:1197; O1-MAD-NEXT: // %bb.0:1198; O1-MAD-NEXT: ld.param.b16 %rs1, [t36_param_0];1199; O1-MAD-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;1200; O1-MAD-NEXT: st.param.b32 [func_retval0], %r1;1201; O1-MAD-NEXT: ret;1202;1203; O0-LABEL: t36(1204; O0: {1205; O0-NEXT: .reg .b16 %rs<3>;1206; O0-NEXT: .reg .b32 %r<3>;1207; O0-EMPTY:1208; O0-NEXT: // %bb.0:1209; O0-NEXT: ld.param.b16 %rs1, [t36_param_0];1210; O0-NEXT: mul.lo.s16 %rs2, %rs1, 3;1211; O0-NEXT: cvt.s32.s16 %r1, %rs2;1212; O0-NEXT: add.s32 %r2, %r1, 5;1213; O0-NEXT: st.param.b32 [func_retval0], %r2;1214; O0-NEXT: ret;1215 %mul = mul nsw i16 %a, 31216 %sext = sext i16 %mul to i321217 %add = add i32 5, %sext1218 ret i32 %add1219}1220 1221define i32 @t37(i16 %a, i16 %b, i32 %c) {1222; O1-NO-MAD-LABEL: t37(1223; O1-NO-MAD: {1224; O1-NO-MAD-NEXT: .reg .b16 %rs<3>;1225; O1-NO-MAD-NEXT: .reg .b32 %r<4>;1226; O1-NO-MAD-EMPTY:1227; O1-NO-MAD-NEXT: // %bb.0:1228; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t37_param_0];1229; O1-NO-MAD-NEXT: ld.param.b16 %rs2, [t37_param_1];1230; O1-NO-MAD-NEXT: mul.wide.u16 %r1, %rs1, %rs2;1231; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t37_param_2];1232; O1-NO-MAD-NEXT: add.s32 %r3, %r2, %r1;1233; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r3;1234; O1-NO-MAD-NEXT: ret;1235;1236; O1-MAD-LABEL: t37(1237; O1-MAD: {1238; O1-MAD-NEXT: .reg .b16 %rs<3>;1239; O1-MAD-NEXT: .reg .b32 %r<3>;1240; O1-MAD-EMPTY:1241; O1-MAD-NEXT: // %bb.0:1242; O1-MAD-NEXT: ld.param.b16 %rs1, [t37_param_0];1243; O1-MAD-NEXT: ld.param.b16 %rs2, [t37_param_1];1244; O1-MAD-NEXT: ld.param.b32 %r1, [t37_param_2];1245; O1-MAD-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;1246; O1-MAD-NEXT: st.param.b32 [func_retval0], %r2;1247; O1-MAD-NEXT: ret;1248;1249; O0-LABEL: t37(1250; O0: {1251; O0-NEXT: .reg .b16 %rs<4>;1252; O0-NEXT: .reg .b32 %r<4>;1253; O0-EMPTY:1254; O0-NEXT: // %bb.0:1255; O0-NEXT: ld.param.b32 %r1, [t37_param_2];1256; O0-NEXT: ld.param.b16 %rs2, [t37_param_1];1257; O0-NEXT: ld.param.b16 %rs1, [t37_param_0];1258; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1259; O0-NEXT: cvt.u32.u16 %r2, %rs3;1260; O0-NEXT: add.s32 %r3, %r1, %r2;1261; O0-NEXT: st.param.b32 [func_retval0], %r3;1262; O0-NEXT: ret;1263 %mul = mul nuw i16 %a, %b1264 %zext = zext i16 %mul to i321265 %add = add i32 %c, %zext1266 ret i32 %add1267}1268 1269define i32 @t38(i16 %a, i16 %b, i32 %c) {1270; O1-NO-MAD-LABEL: t38(1271; O1-NO-MAD: {1272; O1-NO-MAD-NEXT: .reg .b16 %rs<3>;1273; O1-NO-MAD-NEXT: .reg .b32 %r<4>;1274; O1-NO-MAD-EMPTY:1275; O1-NO-MAD-NEXT: // %bb.0:1276; O1-NO-MAD-NEXT: ld.param.b16 %rs1, [t38_param_0];1277; O1-NO-MAD-NEXT: ld.param.b16 %rs2, [t38_param_1];1278; O1-NO-MAD-NEXT: mul.wide.u16 %r1, %rs1, %rs2;1279; O1-NO-MAD-NEXT: ld.param.b32 %r2, [t38_param_2];1280; O1-NO-MAD-NEXT: add.s32 %r3, %r1, %r2;1281; O1-NO-MAD-NEXT: st.param.b32 [func_retval0], %r3;1282; O1-NO-MAD-NEXT: ret;1283;1284; O1-MAD-LABEL: t38(1285; O1-MAD: {1286; O1-MAD-NEXT: .reg .b16 %rs<3>;1287; O1-MAD-NEXT: .reg .b32 %r<3>;1288; O1-MAD-EMPTY:1289; O1-MAD-NEXT: // %bb.0:1290; O1-MAD-NEXT: ld.param.b16 %rs1, [t38_param_0];1291; O1-MAD-NEXT: ld.param.b16 %rs2, [t38_param_1];1292; O1-MAD-NEXT: ld.param.b32 %r1, [t38_param_2];1293; O1-MAD-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;1294; O1-MAD-NEXT: st.param.b32 [func_retval0], %r2;1295; O1-MAD-NEXT: ret;1296;1297; O0-LABEL: t38(1298; O0: {1299; O0-NEXT: .reg .b16 %rs<4>;1300; O0-NEXT: .reg .b32 %r<4>;1301; O0-EMPTY:1302; O0-NEXT: // %bb.0:1303; O0-NEXT: ld.param.b32 %r1, [t38_param_2];1304; O0-NEXT: ld.param.b16 %rs2, [t38_param_1];1305; O0-NEXT: ld.param.b16 %rs1, [t38_param_0];1306; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1307; O0-NEXT: cvt.u32.u16 %r2, %rs3;1308; O0-NEXT: add.s32 %r3, %r2, %r1;1309; O0-NEXT: st.param.b32 [func_retval0], %r3;1310; O0-NEXT: ret;1311 %mul = mul nuw i16 %a, %b1312 %zext = zext i16 %mul to i321313 %add = add i32 %zext, %c1314 ret i32 %add1315}1316 1317define i64 @t39(i16 %a, i16 %b) {1318; O1-LABEL: t39(1319; O1: {1320; O1-NEXT: .reg .b16 %rs<4>;1321; O1-NEXT: .reg .b64 %rd<2>;1322; O1-EMPTY:1323; O1-NEXT: // %bb.0:1324; O1-NEXT: ld.param.b16 %rs1, [t39_param_0];1325; O1-NEXT: ld.param.b16 %rs2, [t39_param_1];1326; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1327; O1-NEXT: cvt.u64.u16 %rd1, %rs3;1328; O1-NEXT: st.param.b64 [func_retval0], %rd1;1329; O1-NEXT: ret;1330;1331; O0-LABEL: t39(1332; O0: {1333; O0-NEXT: .reg .b16 %rs<4>;1334; O0-NEXT: .reg .b64 %rd<2>;1335; O0-EMPTY:1336; O0-NEXT: // %bb.0:1337; O0-NEXT: ld.param.b16 %rs2, [t39_param_1];1338; O0-NEXT: ld.param.b16 %rs1, [t39_param_0];1339; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1340; O0-NEXT: cvt.u64.u16 %rd1, %rs3;1341; O0-NEXT: st.param.b64 [func_retval0], %rd1;1342; O0-NEXT: ret;1343 %mul = mul i16 %a, %b1344 %zext = zext i16 %mul to i641345 ret i64 %zext1346}1347 1348define i64 @t40(i16 %a, i16 %b) {1349; O1-LABEL: t40(1350; O1: {1351; O1-NEXT: .reg .b16 %rs<4>;1352; O1-NEXT: .reg .b64 %rd<2>;1353; O1-EMPTY:1354; O1-NEXT: // %bb.0:1355; O1-NEXT: ld.param.b16 %rs1, [t40_param_0];1356; O1-NEXT: ld.param.b16 %rs2, [t40_param_1];1357; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1358; O1-NEXT: cvt.u64.u16 %rd1, %rs3;1359; O1-NEXT: st.param.b64 [func_retval0], %rd1;1360; O1-NEXT: ret;1361;1362; O0-LABEL: t40(1363; O0: {1364; O0-NEXT: .reg .b16 %rs<4>;1365; O0-NEXT: .reg .b64 %rd<2>;1366; O0-EMPTY:1367; O0-NEXT: // %bb.0:1368; O0-NEXT: ld.param.b16 %rs2, [t40_param_1];1369; O0-NEXT: ld.param.b16 %rs1, [t40_param_0];1370; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1371; O0-NEXT: cvt.u64.u16 %rd1, %rs3;1372; O0-NEXT: st.param.b64 [func_retval0], %rd1;1373; O0-NEXT: ret;1374 %mul = mul nuw i16 %a, %b1375 %zext = zext i16 %mul to i641376 ret i64 %zext1377}1378 1379define i64 @t41(i16 %a, i16 %b) {1380; O1-LABEL: t41(1381; O1: {1382; O1-NEXT: .reg .b16 %rs<4>;1383; O1-NEXT: .reg .b64 %rd<2>;1384; O1-EMPTY:1385; O1-NEXT: // %bb.0:1386; O1-NEXT: ld.param.b16 %rs1, [t41_param_0];1387; O1-NEXT: ld.param.b16 %rs2, [t41_param_1];1388; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1389; O1-NEXT: cvt.s64.s16 %rd1, %rs3;1390; O1-NEXT: st.param.b64 [func_retval0], %rd1;1391; O1-NEXT: ret;1392;1393; O0-LABEL: t41(1394; O0: {1395; O0-NEXT: .reg .b16 %rs<4>;1396; O0-NEXT: .reg .b64 %rd<2>;1397; O0-EMPTY:1398; O0-NEXT: // %bb.0:1399; O0-NEXT: ld.param.b16 %rs2, [t41_param_1];1400; O0-NEXT: ld.param.b16 %rs1, [t41_param_0];1401; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1402; O0-NEXT: cvt.s64.s16 %rd1, %rs3;1403; O0-NEXT: st.param.b64 [func_retval0], %rd1;1404; O0-NEXT: ret;1405 %mul = mul nsw i16 %a, %b1406 %sext = sext i16 %mul to i641407 ret i64 %sext1408}1409 1410define i32 @t42(i16 %a, i16 %b, ptr %ptr) {1411; O1-LABEL: t42(1412; O1: {1413; O1-NEXT: .reg .b16 %rs<4>;1414; O1-NEXT: .reg .b32 %r<2>;1415; O1-NEXT: .reg .b64 %rd<2>;1416; O1-EMPTY:1417; O1-NEXT: // %bb.0:1418; O1-NEXT: ld.param.b16 %rs1, [t42_param_0];1419; O1-NEXT: ld.param.b16 %rs2, [t42_param_1];1420; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1421; O1-NEXT: ld.param.b64 %rd1, [t42_param_2];1422; O1-NEXT: st.b16 [%rd1], %rs3;1423; O1-NEXT: cvt.u32.u16 %r1, %rs3;1424; O1-NEXT: st.param.b32 [func_retval0], %r1;1425; O1-NEXT: ret;1426;1427; O0-LABEL: t42(1428; O0: {1429; O0-NEXT: .reg .b16 %rs<4>;1430; O0-NEXT: .reg .b32 %r<2>;1431; O0-NEXT: .reg .b64 %rd<2>;1432; O0-EMPTY:1433; O0-NEXT: // %bb.0:1434; O0-NEXT: ld.param.b64 %rd1, [t42_param_2];1435; O0-NEXT: ld.param.b16 %rs2, [t42_param_1];1436; O0-NEXT: ld.param.b16 %rs1, [t42_param_0];1437; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1438; O0-NEXT: st.b16 [%rd1], %rs3;1439; O0-NEXT: cvt.u32.u16 %r1, %rs3;1440; O0-NEXT: st.param.b32 [func_retval0], %r1;1441; O0-NEXT: ret;1442 %mul = mul nuw i16 %a, %b1443 store i16 %mul, ptr %ptr1444 %zext = zext i16 %mul to i321445 ret i32 %zext1446}1447 1448define i32 @t43(i16 %a, i16 %b, i32 %c, ptr %ptr) {1449; O1-LABEL: t43(1450; O1: {1451; O1-NEXT: .reg .b16 %rs<4>;1452; O1-NEXT: .reg .b32 %r<4>;1453; O1-NEXT: .reg .b64 %rd<2>;1454; O1-EMPTY:1455; O1-NEXT: // %bb.0:1456; O1-NEXT: ld.param.b16 %rs1, [t43_param_0];1457; O1-NEXT: ld.param.b16 %rs2, [t43_param_1];1458; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1459; O1-NEXT: ld.param.b64 %rd1, [t43_param_3];1460; O1-NEXT: st.b16 [%rd1], %rs3;1461; O1-NEXT: ld.param.b32 %r1, [t43_param_2];1462; O1-NEXT: cvt.u32.u16 %r2, %rs3;1463; O1-NEXT: add.s32 %r3, %r2, %r1;1464; O1-NEXT: st.param.b32 [func_retval0], %r3;1465; O1-NEXT: ret;1466;1467; O0-LABEL: t43(1468; O0: {1469; O0-NEXT: .reg .b16 %rs<4>;1470; O0-NEXT: .reg .b32 %r<4>;1471; O0-NEXT: .reg .b64 %rd<2>;1472; O0-EMPTY:1473; O0-NEXT: // %bb.0:1474; O0-NEXT: ld.param.b64 %rd1, [t43_param_3];1475; O0-NEXT: ld.param.b32 %r1, [t43_param_2];1476; O0-NEXT: ld.param.b16 %rs2, [t43_param_1];1477; O0-NEXT: ld.param.b16 %rs1, [t43_param_0];1478; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;1479; O0-NEXT: st.b16 [%rd1], %rs3;1480; O0-NEXT: cvt.u32.u16 %r2, %rs3;1481; O0-NEXT: add.s32 %r3, %r2, %r1;1482; O0-NEXT: st.param.b32 [func_retval0], %r3;1483; O0-NEXT: ret;1484 %mul = mul nuw i16 %a, %b1485 store i16 %mul, ptr %ptr1486 %zext = zext i16 %mul to i321487 %add = add i32 %zext, %c1488 ret i32 %add1489}1490