339 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | FileCheck %s --check-prefix=OPT3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=NOOPT4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}6 7define i32 @mulwide16(i16 %a, i16 %b) {8; OPT-LABEL: mulwide16(9; OPT: {10; OPT-NEXT: .reg .b16 %rs<3>;11; OPT-NEXT: .reg .b32 %r<2>;12; OPT-EMPTY:13; OPT-NEXT: // %bb.0:14; OPT-NEXT: ld.param.b16 %rs1, [mulwide16_param_0];15; OPT-NEXT: ld.param.b16 %rs2, [mulwide16_param_1];16; OPT-NEXT: mul.wide.s16 %r1, %rs1, %rs2;17; OPT-NEXT: st.param.b32 [func_retval0], %r1;18; OPT-NEXT: ret;19;20; NOOPT-LABEL: mulwide16(21; NOOPT: {22; NOOPT-NEXT: .reg .b16 %rs<3>;23; NOOPT-NEXT: .reg .b32 %r<4>;24; NOOPT-EMPTY:25; NOOPT-NEXT: // %bb.0:26; NOOPT-NEXT: ld.param.b16 %rs2, [mulwide16_param_1];27; NOOPT-NEXT: ld.param.b16 %rs1, [mulwide16_param_0];28; NOOPT-NEXT: cvt.s32.s16 %r1, %rs1;29; NOOPT-NEXT: cvt.s32.s16 %r2, %rs2;30; NOOPT-NEXT: mul.lo.s32 %r3, %r1, %r2;31; NOOPT-NEXT: st.param.b32 [func_retval0], %r3;32; NOOPT-NEXT: ret;33 %val0 = sext i16 %a to i3234 %val1 = sext i16 %b to i3235 %val2 = mul i32 %val0, %val136 ret i32 %val237}38 39define i32 @mulwideu16(i16 %a, i16 %b) {40; OPT-LABEL: mulwideu16(41; OPT: {42; OPT-NEXT: .reg .b16 %rs<3>;43; OPT-NEXT: .reg .b32 %r<2>;44; OPT-EMPTY:45; OPT-NEXT: // %bb.0:46; OPT-NEXT: ld.param.b16 %rs1, [mulwideu16_param_0];47; OPT-NEXT: ld.param.b16 %rs2, [mulwideu16_param_1];48; OPT-NEXT: mul.wide.u16 %r1, %rs1, %rs2;49; OPT-NEXT: st.param.b32 [func_retval0], %r1;50; OPT-NEXT: ret;51;52; NOOPT-LABEL: mulwideu16(53; NOOPT: {54; NOOPT-NEXT: .reg .b16 %rs<3>;55; NOOPT-NEXT: .reg .b32 %r<4>;56; NOOPT-EMPTY:57; NOOPT-NEXT: // %bb.0:58; NOOPT-NEXT: ld.param.b16 %rs2, [mulwideu16_param_1];59; NOOPT-NEXT: ld.param.b16 %rs1, [mulwideu16_param_0];60; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;61; NOOPT-NEXT: cvt.u32.u16 %r2, %rs2;62; NOOPT-NEXT: mul.lo.s32 %r3, %r1, %r2;63; NOOPT-NEXT: st.param.b32 [func_retval0], %r3;64; NOOPT-NEXT: ret;65 %val0 = zext i16 %a to i3266 %val1 = zext i16 %b to i3267 %val2 = mul i32 %val0, %val168 ret i32 %val269}70 71define i32 @mulwide8(i8 %a, i8 %b) {72; OPT-LABEL: mulwide8(73; OPT: {74; OPT-NEXT: .reg .b16 %rs<3>;75; OPT-NEXT: .reg .b32 %r<2>;76; OPT-EMPTY:77; OPT-NEXT: // %bb.0:78; OPT-NEXT: ld.param.s8 %rs1, [mulwide8_param_0];79; OPT-NEXT: ld.param.s8 %rs2, [mulwide8_param_1];80; OPT-NEXT: mul.wide.s16 %r1, %rs1, %rs2;81; OPT-NEXT: st.param.b32 [func_retval0], %r1;82; OPT-NEXT: ret;83;84; NOOPT-LABEL: mulwide8(85; NOOPT: {86; NOOPT-NEXT: .reg .b16 %rs<3>;87; NOOPT-NEXT: .reg .b32 %r<6>;88; NOOPT-EMPTY:89; NOOPT-NEXT: // %bb.0:90; NOOPT-NEXT: ld.param.b8 %rs2, [mulwide8_param_1];91; NOOPT-NEXT: ld.param.b8 %rs1, [mulwide8_param_0];92; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;93; NOOPT-NEXT: cvt.s32.s8 %r2, %r1;94; NOOPT-NEXT: cvt.u32.u16 %r3, %rs2;95; NOOPT-NEXT: cvt.s32.s8 %r4, %r3;96; NOOPT-NEXT: mul.lo.s32 %r5, %r2, %r4;97; NOOPT-NEXT: st.param.b32 [func_retval0], %r5;98; NOOPT-NEXT: ret;99 %val0 = sext i8 %a to i32100 %val1 = sext i8 %b to i32101 %val2 = mul i32 %val0, %val1102 ret i32 %val2103}104 105define i32 @mulwideu8(i8 %a, i8 %b) {106; OPT-LABEL: mulwideu8(107; OPT: {108; OPT-NEXT: .reg .b16 %rs<3>;109; OPT-NEXT: .reg .b32 %r<2>;110; OPT-EMPTY:111; OPT-NEXT: // %bb.0:112; OPT-NEXT: ld.param.b8 %rs1, [mulwideu8_param_0];113; OPT-NEXT: ld.param.b8 %rs2, [mulwideu8_param_1];114; OPT-NEXT: mul.wide.u16 %r1, %rs1, %rs2;115; OPT-NEXT: st.param.b32 [func_retval0], %r1;116; OPT-NEXT: ret;117;118; NOOPT-LABEL: mulwideu8(119; NOOPT: {120; NOOPT-NEXT: .reg .b16 %rs<3>;121; NOOPT-NEXT: .reg .b32 %r<4>;122; NOOPT-EMPTY:123; NOOPT-NEXT: // %bb.0:124; NOOPT-NEXT: ld.param.b8 %rs2, [mulwideu8_param_1];125; NOOPT-NEXT: ld.param.b8 %rs1, [mulwideu8_param_0];126; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;127; NOOPT-NEXT: cvt.u32.u16 %r2, %rs2;128; NOOPT-NEXT: mul.lo.s32 %r3, %r1, %r2;129; NOOPT-NEXT: st.param.b32 [func_retval0], %r3;130; NOOPT-NEXT: ret;131 %val0 = zext i8 %a to i32132 %val1 = zext i8 %b to i32133 %val2 = mul i32 %val0, %val1134 ret i32 %val2135}136 137define i64 @mulwide32(i32 %a, i32 %b) {138; OPT-LABEL: mulwide32(139; OPT: {140; OPT-NEXT: .reg .b32 %r<3>;141; OPT-NEXT: .reg .b64 %rd<2>;142; OPT-EMPTY:143; OPT-NEXT: // %bb.0:144; OPT-NEXT: ld.param.b32 %r1, [mulwide32_param_0];145; OPT-NEXT: ld.param.b32 %r2, [mulwide32_param_1];146; OPT-NEXT: mul.wide.s32 %rd1, %r1, %r2;147; OPT-NEXT: st.param.b64 [func_retval0], %rd1;148; OPT-NEXT: ret;149;150; NOOPT-LABEL: mulwide32(151; NOOPT: {152; NOOPT-NEXT: .reg .b32 %r<3>;153; NOOPT-NEXT: .reg .b64 %rd<4>;154; NOOPT-EMPTY:155; NOOPT-NEXT: // %bb.0:156; NOOPT-NEXT: ld.param.b32 %r2, [mulwide32_param_1];157; NOOPT-NEXT: ld.param.b32 %r1, [mulwide32_param_0];158; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;159; NOOPT-NEXT: cvt.s64.s32 %rd2, %r2;160; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;161; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;162; NOOPT-NEXT: ret;163 %val0 = sext i32 %a to i64164 %val1 = sext i32 %b to i64165 %val2 = mul i64 %val0, %val1166 ret i64 %val2167}168 169define i64 @mulwideu32(i32 %a, i32 %b) {170; OPT-LABEL: mulwideu32(171; OPT: {172; OPT-NEXT: .reg .b32 %r<3>;173; OPT-NEXT: .reg .b64 %rd<2>;174; OPT-EMPTY:175; OPT-NEXT: // %bb.0:176; OPT-NEXT: ld.param.b32 %r1, [mulwideu32_param_0];177; OPT-NEXT: ld.param.b32 %r2, [mulwideu32_param_1];178; OPT-NEXT: mul.wide.u32 %rd1, %r1, %r2;179; OPT-NEXT: st.param.b64 [func_retval0], %rd1;180; OPT-NEXT: ret;181;182; NOOPT-LABEL: mulwideu32(183; NOOPT: {184; NOOPT-NEXT: .reg .b32 %r<3>;185; NOOPT-NEXT: .reg .b64 %rd<4>;186; NOOPT-EMPTY:187; NOOPT-NEXT: // %bb.0:188; NOOPT-NEXT: ld.param.b32 %r2, [mulwideu32_param_1];189; NOOPT-NEXT: ld.param.b32 %r1, [mulwideu32_param_0];190; NOOPT-NEXT: cvt.u64.u32 %rd1, %r1;191; NOOPT-NEXT: cvt.u64.u32 %rd2, %r2;192; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;193; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;194; NOOPT-NEXT: ret;195 %val0 = zext i32 %a to i64196 %val1 = zext i32 %b to i64197 %val2 = mul i64 %val0, %val1198 ret i64 %val2199}200 201define i64 @mulwideu7(i7 %a, i7 %b) {202; OPT-LABEL: mulwideu7(203; OPT: {204; OPT-NEXT: .reg .b32 %r<5>;205; OPT-NEXT: .reg .b64 %rd<2>;206; OPT-EMPTY:207; OPT-NEXT: // %bb.0:208; OPT-NEXT: ld.param.b8 %r1, [mulwideu7_param_1];209; OPT-NEXT: and.b32 %r2, %r1, 127;210; OPT-NEXT: ld.param.b8 %r3, [mulwideu7_param_0];211; OPT-NEXT: and.b32 %r4, %r3, 127;212; OPT-NEXT: mul.wide.u32 %rd1, %r4, %r2;213; OPT-NEXT: st.param.b64 [func_retval0], %rd1;214; OPT-NEXT: ret;215;216; NOOPT-LABEL: mulwideu7(217; NOOPT: {218; NOOPT-NEXT: .reg .b16 %rs<9>;219; NOOPT-NEXT: .reg .b64 %rd<6>;220; NOOPT-EMPTY:221; NOOPT-NEXT: // %bb.0:222; NOOPT-NEXT: ld.param.b8 %rs3, [mulwideu7_param_0+1];223; NOOPT-NEXT: shl.b16 %rs4, %rs3, 8;224; NOOPT-NEXT: ld.param.b8 %rs5, [mulwideu7_param_0];225; NOOPT-NEXT: or.b16 %rs1, %rs4, %rs5;226; NOOPT-NEXT: ld.param.b8 %rs6, [mulwideu7_param_1+1];227; NOOPT-NEXT: shl.b16 %rs7, %rs6, 8;228; NOOPT-NEXT: ld.param.b8 %rs8, [mulwideu7_param_1];229; NOOPT-NEXT: or.b16 %rs2, %rs7, %rs8;230; NOOPT-NEXT: cvt.u64.u16 %rd1, %rs5;231; NOOPT-NEXT: and.b64 %rd2, %rd1, 127;232; NOOPT-NEXT: cvt.u64.u16 %rd3, %rs8;233; NOOPT-NEXT: and.b64 %rd4, %rd3, 127;234; NOOPT-NEXT: mul.lo.s64 %rd5, %rd2, %rd4;235; NOOPT-NEXT: st.param.b64 [func_retval0], %rd5;236; NOOPT-NEXT: ret;237 %val0 = zext i7 %a to i64238 %val1 = zext i7 %b to i64239 %val2 = mul i64 %val0, %val1240 ret i64 %val2241}242 243define i64 @mulwides7(i7 %a, i7 %b) {244; OPT-LABEL: mulwides7(245; OPT: {246; OPT-NEXT: .reg .b32 %r<5>;247; OPT-NEXT: .reg .b64 %rd<2>;248; OPT-EMPTY:249; OPT-NEXT: // %bb.0:250; OPT-NEXT: ld.param.b8 %r1, [mulwides7_param_0];251; OPT-NEXT: ld.param.b8 %r2, [mulwides7_param_1];252; OPT-NEXT: bfe.s32 %r3, %r2, 0, 7;253; OPT-NEXT: bfe.s32 %r4, %r1, 0, 7;254; OPT-NEXT: mul.wide.s32 %rd1, %r4, %r3;255; OPT-NEXT: st.param.b64 [func_retval0], %rd1;256; OPT-NEXT: ret;257;258; NOOPT-LABEL: mulwides7(259; NOOPT: {260; NOOPT-NEXT: .reg .b16 %rs<9>;261; NOOPT-NEXT: .reg .b64 %rd<6>;262; NOOPT-EMPTY:263; NOOPT-NEXT: // %bb.0:264; NOOPT-NEXT: ld.param.b8 %rs3, [mulwides7_param_0+1];265; NOOPT-NEXT: shl.b16 %rs4, %rs3, 8;266; NOOPT-NEXT: ld.param.b8 %rs5, [mulwides7_param_0];267; NOOPT-NEXT: or.b16 %rs1, %rs4, %rs5;268; NOOPT-NEXT: ld.param.b8 %rs6, [mulwides7_param_1];269; NOOPT-NEXT: cvt.u64.u16 %rd1, %rs6;270; NOOPT-NEXT: cvt.u64.u16 %rd2, %rs5;271; NOOPT-NEXT: ld.param.b8 %rs7, [mulwides7_param_1+1];272; NOOPT-NEXT: shl.b16 %rs8, %rs7, 8;273; NOOPT-NEXT: or.b16 %rs2, %rs8, %rs6;274; NOOPT-NEXT: bfe.s64 %rd3, %rd2, 0, 7;275; NOOPT-NEXT: bfe.s64 %rd4, %rd1, 0, 7;276; NOOPT-NEXT: mul.lo.s64 %rd5, %rd3, %rd4;277; NOOPT-NEXT: st.param.b64 [func_retval0], %rd5;278; NOOPT-NEXT: ret;279 %val0 = sext i7 %a to i64280 %val1 = sext i7 %b to i64281 %val2 = mul i64 %val0, %val1282 ret i64 %val2283}284 285define i64 @shl30(i32 %a) {286; OPT-LABEL: shl30(287; OPT: {288; OPT-NEXT: .reg .b32 %r<2>;289; OPT-NEXT: .reg .b64 %rd<2>;290; OPT-EMPTY:291; OPT-NEXT: // %bb.0:292; OPT-NEXT: ld.param.b32 %r1, [shl30_param_0];293; OPT-NEXT: mul.wide.s32 %rd1, %r1, 1073741824;294; OPT-NEXT: st.param.b64 [func_retval0], %rd1;295; OPT-NEXT: ret;296;297; NOOPT-LABEL: shl30(298; NOOPT: {299; NOOPT-NEXT: .reg .b32 %r<2>;300; NOOPT-NEXT: .reg .b64 %rd<3>;301; NOOPT-EMPTY:302; NOOPT-NEXT: // %bb.0:303; NOOPT-NEXT: ld.param.b32 %r1, [shl30_param_0];304; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;305; NOOPT-NEXT: shl.b64 %rd2, %rd1, 30;306; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;307; NOOPT-NEXT: ret;308 %conv = sext i32 %a to i64309 %shl = shl i64 %conv, 30310 ret i64 %shl311}312 313define i64 @shl31(i32 %a) {314; OPT-LABEL: shl31(315; OPT: {316; OPT-NEXT: .reg .b64 %rd<3>;317; OPT-EMPTY:318; OPT-NEXT: // %bb.0:319; OPT-NEXT: ld.param.s32 %rd1, [shl31_param_0];320; OPT-NEXT: shl.b64 %rd2, %rd1, 31;321; OPT-NEXT: st.param.b64 [func_retval0], %rd2;322; OPT-NEXT: ret;323;324; NOOPT-LABEL: shl31(325; NOOPT: {326; NOOPT-NEXT: .reg .b32 %r<2>;327; NOOPT-NEXT: .reg .b64 %rd<3>;328; NOOPT-EMPTY:329; NOOPT-NEXT: // %bb.0:330; NOOPT-NEXT: ld.param.b32 %r1, [shl31_param_0];331; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;332; NOOPT-NEXT: shl.b64 %rd2, %rd1, 31;333; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;334; NOOPT-NEXT: ret;335 %conv = sext i32 %a to i64336 %shl = shl i64 %conv, 31337 ret i64 %shl338}339