748 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc < %s --mtriple=nvptx64 -mcpu=sm_20 | FileCheck --check-prefix=SM20 %s3; RUN: llc < %s --mtriple=nvptx64 -mcpu=sm_35 | FileCheck --check-prefix=SM35 %s4; RUN: %if ptxas %{ llc < %s --mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s --mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}6 7 8declare i32 @llvm.nvvm.rotate.b32(i32, i32)9declare i64 @llvm.nvvm.rotate.b64(i64, i32)10declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)11 12declare i64 @llvm.fshl.i64(i64, i64, i64)13declare i64 @llvm.fshr.i64(i64, i64, i64)14declare i32 @llvm.fshl.i32(i32, i32, i32)15declare i32 @llvm.fshr.i32(i32, i32, i32)16 17 18define i32 @rotate32(i32 %a, i32 %b) {19; SM20-LABEL: rotate32(20; SM20: {21; SM20-NEXT: .reg .b32 %r<9>;22; SM20-EMPTY:23; SM20-NEXT: // %bb.0:24; SM20-NEXT: ld.param.b32 %r1, [rotate32_param_0];25; SM20-NEXT: ld.param.b32 %r2, [rotate32_param_1];26; SM20-NEXT: and.b32 %r3, %r2, 31;27; SM20-NEXT: shl.b32 %r4, %r1, %r3;28; SM20-NEXT: neg.s32 %r5, %r2;29; SM20-NEXT: and.b32 %r6, %r5, 31;30; SM20-NEXT: shr.u32 %r7, %r1, %r6;31; SM20-NEXT: or.b32 %r8, %r4, %r7;32; SM20-NEXT: st.param.b32 [func_retval0], %r8;33; SM20-NEXT: ret;34;35; SM35-LABEL: rotate32(36; SM35: {37; SM35-NEXT: .reg .b32 %r<4>;38; SM35-EMPTY:39; SM35-NEXT: // %bb.0:40; SM35-NEXT: ld.param.b32 %r1, [rotate32_param_0];41; SM35-NEXT: ld.param.b32 %r2, [rotate32_param_1];42; SM35-NEXT: shf.l.wrap.b32 %r3, %r1, %r1, %r2;43; SM35-NEXT: st.param.b32 [func_retval0], %r3;44; SM35-NEXT: ret;45 %val = tail call i32 @llvm.nvvm.rotate.b32(i32 %a, i32 %b)46 ret i32 %val47}48 49define i64 @rotate64(i64 %a, i32 %b) {50; SM20-LABEL: rotate64(51; SM20: {52; SM20-NEXT: .reg .b32 %r<5>;53; SM20-NEXT: .reg .b64 %rd<5>;54; SM20-EMPTY:55; SM20-NEXT: // %bb.0:56; SM20-NEXT: ld.param.b64 %rd1, [rotate64_param_0];57; SM20-NEXT: ld.param.b32 %r1, [rotate64_param_1];58; SM20-NEXT: and.b32 %r2, %r1, 63;59; SM20-NEXT: shl.b64 %rd2, %rd1, %r2;60; SM20-NEXT: neg.s32 %r3, %r1;61; SM20-NEXT: and.b32 %r4, %r3, 63;62; SM20-NEXT: shr.u64 %rd3, %rd1, %r4;63; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;64; SM20-NEXT: st.param.b64 [func_retval0], %rd4;65; SM20-NEXT: ret;66;67; SM35-LABEL: rotate64(68; SM35: {69; SM35-NEXT: .reg .b32 %r<5>;70; SM35-NEXT: .reg .b64 %rd<5>;71; SM35-EMPTY:72; SM35-NEXT: // %bb.0:73; SM35-NEXT: ld.param.b64 %rd1, [rotate64_param_0];74; SM35-NEXT: ld.param.b32 %r1, [rotate64_param_1];75; SM35-NEXT: and.b32 %r2, %r1, 63;76; SM35-NEXT: shl.b64 %rd2, %rd1, %r2;77; SM35-NEXT: neg.s32 %r3, %r1;78; SM35-NEXT: and.b32 %r4, %r3, 63;79; SM35-NEXT: shr.u64 %rd3, %rd1, %r4;80; SM35-NEXT: or.b64 %rd4, %rd2, %rd3;81; SM35-NEXT: st.param.b64 [func_retval0], %rd4;82; SM35-NEXT: ret;83 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 %b)84 ret i64 %val85}86 87define i64 @rotateright64(i64 %a, i32 %b) {88; SM20-LABEL: rotateright64(89; SM20: {90; SM20-NEXT: .reg .b32 %r<5>;91; SM20-NEXT: .reg .b64 %rd<5>;92; SM20-EMPTY:93; SM20-NEXT: // %bb.0:94; SM20-NEXT: ld.param.b64 %rd1, [rotateright64_param_0];95; SM20-NEXT: ld.param.b32 %r1, [rotateright64_param_1];96; SM20-NEXT: and.b32 %r2, %r1, 63;97; SM20-NEXT: shr.u64 %rd2, %rd1, %r2;98; SM20-NEXT: neg.s32 %r3, %r1;99; SM20-NEXT: and.b32 %r4, %r3, 63;100; SM20-NEXT: shl.b64 %rd3, %rd1, %r4;101; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;102; SM20-NEXT: st.param.b64 [func_retval0], %rd4;103; SM20-NEXT: ret;104;105; SM35-LABEL: rotateright64(106; SM35: {107; SM35-NEXT: .reg .b32 %r<5>;108; SM35-NEXT: .reg .b64 %rd<5>;109; SM35-EMPTY:110; SM35-NEXT: // %bb.0:111; SM35-NEXT: ld.param.b64 %rd1, [rotateright64_param_0];112; SM35-NEXT: ld.param.b32 %r1, [rotateright64_param_1];113; SM35-NEXT: and.b32 %r2, %r1, 63;114; SM35-NEXT: shr.u64 %rd2, %rd1, %r2;115; SM35-NEXT: neg.s32 %r3, %r1;116; SM35-NEXT: and.b32 %r4, %r3, 63;117; SM35-NEXT: shl.b64 %rd3, %rd1, %r4;118; SM35-NEXT: or.b64 %rd4, %rd2, %rd3;119; SM35-NEXT: st.param.b64 [func_retval0], %rd4;120; SM35-NEXT: ret;121 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 %b)122 ret i64 %val123}124 125define i32 @rotl0(i32 %x) {126; SM20-LABEL: rotl0(127; SM20: {128; SM20-NEXT: .reg .b32 %r<5>;129; SM20-EMPTY:130; SM20-NEXT: // %bb.0:131; SM20-NEXT: ld.param.b32 %r1, [rotl0_param_0];132; SM20-NEXT: shr.u32 %r2, %r1, 24;133; SM20-NEXT: shl.b32 %r3, %r1, 8;134; SM20-NEXT: or.b32 %r4, %r3, %r2;135; SM20-NEXT: st.param.b32 [func_retval0], %r4;136; SM20-NEXT: ret;137;138; SM35-LABEL: rotl0(139; SM35: {140; SM35-NEXT: .reg .b32 %r<3>;141; SM35-EMPTY:142; SM35-NEXT: // %bb.0:143; SM35-NEXT: ld.param.b32 %r1, [rotl0_param_0];144; SM35-NEXT: shf.l.wrap.b32 %r2, %r1, %r1, 8;145; SM35-NEXT: st.param.b32 [func_retval0], %r2;146; SM35-NEXT: ret;147 %t0 = shl i32 %x, 8148 %t1 = lshr i32 %x, 24149 %t2 = or i32 %t0, %t1150 ret i32 %t2151}152 153define i64 @rotl64(i64 %a, i64 %n) {154; SM20-LABEL: rotl64(155; SM20: {156; SM20-NEXT: .reg .b32 %r<5>;157; SM20-NEXT: .reg .b64 %rd<5>;158; SM20-EMPTY:159; SM20-NEXT: // %bb.0:160; SM20-NEXT: ld.param.b64 %rd1, [rotl64_param_0];161; SM20-NEXT: ld.param.b32 %r1, [rotl64_param_1];162; SM20-NEXT: and.b32 %r2, %r1, 63;163; SM20-NEXT: shl.b64 %rd2, %rd1, %r2;164; SM20-NEXT: neg.s32 %r3, %r1;165; SM20-NEXT: and.b32 %r4, %r3, 63;166; SM20-NEXT: shr.u64 %rd3, %rd1, %r4;167; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;168; SM20-NEXT: st.param.b64 [func_retval0], %rd4;169; SM20-NEXT: ret;170;171; SM35-LABEL: rotl64(172; SM35: {173; SM35-NEXT: .reg .b32 %r<5>;174; SM35-NEXT: .reg .b64 %rd<5>;175; SM35-EMPTY:176; SM35-NEXT: // %bb.0:177; SM35-NEXT: ld.param.b64 %rd1, [rotl64_param_0];178; SM35-NEXT: ld.param.b32 %r1, [rotl64_param_1];179; SM35-NEXT: and.b32 %r2, %r1, 63;180; SM35-NEXT: shl.b64 %rd2, %rd1, %r2;181; SM35-NEXT: neg.s32 %r3, %r1;182; SM35-NEXT: and.b32 %r4, %r3, 63;183; SM35-NEXT: shr.u64 %rd3, %rd1, %r4;184; SM35-NEXT: or.b64 %rd4, %rd2, %rd3;185; SM35-NEXT: st.param.b64 [func_retval0], %rd4;186; SM35-NEXT: ret;187 %val = tail call i64 @llvm.fshl.i64(i64 %a, i64 %a, i64 %n)188 ret i64 %val189}190 191define i64 @rotl64_low_imm(i64 %a) {192; SM20-LABEL: rotl64_low_imm(193; SM20: {194; SM20-NEXT: .reg .b64 %rd<5>;195; SM20-EMPTY:196; SM20-NEXT: // %bb.0:197; SM20-NEXT: ld.param.b64 %rd1, [rotl64_low_imm_param_0];198; SM20-NEXT: shr.u64 %rd2, %rd1, 62;199; SM20-NEXT: shl.b64 %rd3, %rd1, 2;200; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;201; SM20-NEXT: st.param.b64 [func_retval0], %rd4;202; SM20-NEXT: ret;203;204; SM35-LABEL: rotl64_low_imm(205; SM35: {206; SM35-NEXT: .reg .b32 %r<5>;207; SM35-NEXT: .reg .b64 %rd<3>;208; SM35-EMPTY:209; SM35-NEXT: // %bb.0:210; SM35-NEXT: ld.param.b64 %rd1, [rotl64_low_imm_param_0];211; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;212; SM35-NEXT: shf.l.wrap.b32 %r3, %r1, %r2, 2;213; SM35-NEXT: shf.l.wrap.b32 %r4, %r2, %r1, 2;214; SM35-NEXT: mov.b64 %rd2, {%r4, %r3};215; SM35-NEXT: st.param.b64 [func_retval0], %rd2;216; SM35-NEXT: ret;217 %val = tail call i64 @llvm.fshl.i64(i64 %a, i64 %a, i64 66)218 ret i64 %val219}220 221define i64 @rotl64_high_imm(i64 %a) {222; SM20-LABEL: rotl64_high_imm(223; SM20: {224; SM20-NEXT: .reg .b64 %rd<5>;225; SM20-EMPTY:226; SM20-NEXT: // %bb.0:227; SM20-NEXT: ld.param.b64 %rd1, [rotl64_high_imm_param_0];228; SM20-NEXT: shr.u64 %rd2, %rd1, 1;229; SM20-NEXT: shl.b64 %rd3, %rd1, 63;230; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;231; SM20-NEXT: st.param.b64 [func_retval0], %rd4;232; SM20-NEXT: ret;233;234; SM35-LABEL: rotl64_high_imm(235; SM35: {236; SM35-NEXT: .reg .b32 %r<5>;237; SM35-NEXT: .reg .b64 %rd<3>;238; SM35-EMPTY:239; SM35-NEXT: // %bb.0:240; SM35-NEXT: ld.param.b64 %rd1, [rotl64_high_imm_param_0];241; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;242; SM35-NEXT: shf.l.wrap.b32 %r3, %r2, %r1, 31;243; SM35-NEXT: shf.l.wrap.b32 %r4, %r1, %r2, 31;244; SM35-NEXT: mov.b64 %rd2, {%r4, %r3};245; SM35-NEXT: st.param.b64 [func_retval0], %rd2;246; SM35-NEXT: ret;247 %val = tail call i64 @llvm.fshl.i64(i64 %a, i64 %a, i64 63)248 ret i64 %val249}250 251define i64 @rotl64_32_imm(i64 %a) {252; SM20-LABEL: rotl64_32_imm(253; SM20: {254; SM20-NEXT: .reg .b64 %rd<5>;255; SM20-EMPTY:256; SM20-NEXT: // %bb.0:257; SM20-NEXT: ld.param.b64 %rd1, [rotl64_32_imm_param_0];258; SM20-NEXT: shr.u64 %rd2, %rd1, 32;259; SM20-NEXT: shl.b64 %rd3, %rd1, 32;260; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;261; SM20-NEXT: st.param.b64 [func_retval0], %rd4;262; SM20-NEXT: ret;263;264; SM35-LABEL: rotl64_32_imm(265; SM35: {266; SM35-NEXT: .reg .b32 %r<3>;267; SM35-NEXT: .reg .b64 %rd<3>;268; SM35-EMPTY:269; SM35-NEXT: // %bb.0:270; SM35-NEXT: ld.param.b64 %rd1, [rotl64_32_imm_param_0];271; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;272; SM35-NEXT: mov.b64 %rd2, {%r2, %r1};273; SM35-NEXT: st.param.b64 [func_retval0], %rd2;274; SM35-NEXT: ret;275 %val = tail call i64 @llvm.fshl.i64(i64 %a, i64 %a, i64 32)276 ret i64 %val277}278 279define i64 @rotr64(i64 %a, i64 %n) {280; SM20-LABEL: rotr64(281; SM20: {282; SM20-NEXT: .reg .b32 %r<5>;283; SM20-NEXT: .reg .b64 %rd<5>;284; SM20-EMPTY:285; SM20-NEXT: // %bb.0:286; SM20-NEXT: ld.param.b64 %rd1, [rotr64_param_0];287; SM20-NEXT: ld.param.b32 %r1, [rotr64_param_1];288; SM20-NEXT: and.b32 %r2, %r1, 63;289; SM20-NEXT: shr.u64 %rd2, %rd1, %r2;290; SM20-NEXT: neg.s32 %r3, %r1;291; SM20-NEXT: and.b32 %r4, %r3, 63;292; SM20-NEXT: shl.b64 %rd3, %rd1, %r4;293; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;294; SM20-NEXT: st.param.b64 [func_retval0], %rd4;295; SM20-NEXT: ret;296;297; SM35-LABEL: rotr64(298; SM35: {299; SM35-NEXT: .reg .b32 %r<5>;300; SM35-NEXT: .reg .b64 %rd<5>;301; SM35-EMPTY:302; SM35-NEXT: // %bb.0:303; SM35-NEXT: ld.param.b64 %rd1, [rotr64_param_0];304; SM35-NEXT: ld.param.b32 %r1, [rotr64_param_1];305; SM35-NEXT: and.b32 %r2, %r1, 63;306; SM35-NEXT: shr.u64 %rd2, %rd1, %r2;307; SM35-NEXT: neg.s32 %r3, %r1;308; SM35-NEXT: and.b32 %r4, %r3, 63;309; SM35-NEXT: shl.b64 %rd3, %rd1, %r4;310; SM35-NEXT: or.b64 %rd4, %rd2, %rd3;311; SM35-NEXT: st.param.b64 [func_retval0], %rd4;312; SM35-NEXT: ret;313 %val = tail call i64 @llvm.fshr.i64(i64 %a, i64 %a, i64 %n)314 ret i64 %val315}316 317define i64 @rotr64_low_imm(i64 %a) {318; SM20-LABEL: rotr64_low_imm(319; SM20: {320; SM20-NEXT: .reg .b64 %rd<5>;321; SM20-EMPTY:322; SM20-NEXT: // %bb.0:323; SM20-NEXT: ld.param.b64 %rd1, [rotr64_low_imm_param_0];324; SM20-NEXT: shl.b64 %rd2, %rd1, 52;325; SM20-NEXT: shr.u64 %rd3, %rd1, 12;326; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;327; SM20-NEXT: st.param.b64 [func_retval0], %rd4;328; SM20-NEXT: ret;329;330; SM35-LABEL: rotr64_low_imm(331; SM35: {332; SM35-NEXT: .reg .b32 %r<5>;333; SM35-NEXT: .reg .b64 %rd<3>;334; SM35-EMPTY:335; SM35-NEXT: // %bb.0:336; SM35-NEXT: ld.param.b64 %rd1, [rotr64_low_imm_param_0];337; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;338; SM35-NEXT: shf.r.wrap.b32 %r3, %r2, %r1, 12;339; SM35-NEXT: shf.r.wrap.b32 %r4, %r1, %r2, 12;340; SM35-NEXT: mov.b64 %rd2, {%r4, %r3};341; SM35-NEXT: st.param.b64 [func_retval0], %rd2;342; SM35-NEXT: ret;343 %val = tail call i64 @llvm.fshr.i64(i64 %a, i64 %a, i64 12)344 ret i64 %val345}346 347define i64 @rotr64_high_imm(i64 %a) {348; SM20-LABEL: rotr64_high_imm(349; SM20: {350; SM20-NEXT: .reg .b64 %rd<5>;351; SM20-EMPTY:352; SM20-NEXT: // %bb.0:353; SM20-NEXT: ld.param.b64 %rd1, [rotr64_high_imm_param_0];354; SM20-NEXT: shl.b64 %rd2, %rd1, 21;355; SM20-NEXT: shr.u64 %rd3, %rd1, 43;356; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;357; SM20-NEXT: st.param.b64 [func_retval0], %rd4;358; SM20-NEXT: ret;359;360; SM35-LABEL: rotr64_high_imm(361; SM35: {362; SM35-NEXT: .reg .b32 %r<5>;363; SM35-NEXT: .reg .b64 %rd<3>;364; SM35-EMPTY:365; SM35-NEXT: // %bb.0:366; SM35-NEXT: ld.param.b64 %rd1, [rotr64_high_imm_param_0];367; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;368; SM35-NEXT: shf.r.wrap.b32 %r3, %r1, %r2, 11;369; SM35-NEXT: shf.r.wrap.b32 %r4, %r2, %r1, 11;370; SM35-NEXT: mov.b64 %rd2, {%r4, %r3};371; SM35-NEXT: st.param.b64 [func_retval0], %rd2;372; SM35-NEXT: ret;373 %val = tail call i64 @llvm.fshr.i64(i64 %a, i64 %a, i64 43)374 ret i64 %val375}376 377define i64 @rotr64_32_imm(i64 %a) {378; SM20-LABEL: rotr64_32_imm(379; SM20: {380; SM20-NEXT: .reg .b64 %rd<5>;381; SM20-EMPTY:382; SM20-NEXT: // %bb.0:383; SM20-NEXT: ld.param.b64 %rd1, [rotr64_32_imm_param_0];384; SM20-NEXT: shl.b64 %rd2, %rd1, 32;385; SM20-NEXT: shr.u64 %rd3, %rd1, 32;386; SM20-NEXT: or.b64 %rd4, %rd3, %rd2;387; SM20-NEXT: st.param.b64 [func_retval0], %rd4;388; SM20-NEXT: ret;389;390; SM35-LABEL: rotr64_32_imm(391; SM35: {392; SM35-NEXT: .reg .b32 %r<3>;393; SM35-NEXT: .reg .b64 %rd<3>;394; SM35-EMPTY:395; SM35-NEXT: // %bb.0:396; SM35-NEXT: ld.param.b64 %rd1, [rotr64_32_imm_param_0];397; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;398; SM35-NEXT: mov.b64 %rd2, {%r2, %r1};399; SM35-NEXT: st.param.b64 [func_retval0], %rd2;400; SM35-NEXT: ret;401 %val = tail call i64 @llvm.fshr.i64(i64 %a, i64 %a, i64 32)402 ret i64 %val403}404 405define i32 @funnel_shift_right_32(i32 %a, i32 %b, i32 %c) {406; SM20-LABEL: funnel_shift_right_32(407; SM20: {408; SM20-NEXT: .reg .b32 %r<11>;409; SM20-EMPTY:410; SM20-NEXT: // %bb.0:411; SM20-NEXT: ld.param.b32 %r1, [funnel_shift_right_32_param_0];412; SM20-NEXT: ld.param.b32 %r2, [funnel_shift_right_32_param_2];413; SM20-NEXT: and.b32 %r3, %r2, 31;414; SM20-NEXT: ld.param.b32 %r4, [funnel_shift_right_32_param_1];415; SM20-NEXT: shr.u32 %r5, %r4, %r3;416; SM20-NEXT: shl.b32 %r6, %r1, 1;417; SM20-NEXT: not.b32 %r7, %r2;418; SM20-NEXT: and.b32 %r8, %r7, 31;419; SM20-NEXT: shl.b32 %r9, %r6, %r8;420; SM20-NEXT: or.b32 %r10, %r9, %r5;421; SM20-NEXT: st.param.b32 [func_retval0], %r10;422; SM20-NEXT: ret;423;424; SM35-LABEL: funnel_shift_right_32(425; SM35: {426; SM35-NEXT: .reg .b32 %r<5>;427; SM35-EMPTY:428; SM35-NEXT: // %bb.0:429; SM35-NEXT: ld.param.b32 %r1, [funnel_shift_right_32_param_0];430; SM35-NEXT: ld.param.b32 %r2, [funnel_shift_right_32_param_1];431; SM35-NEXT: ld.param.b32 %r3, [funnel_shift_right_32_param_2];432; SM35-NEXT: shf.r.wrap.b32 %r4, %r2, %r1, %r3;433; SM35-NEXT: st.param.b32 [func_retval0], %r4;434; SM35-NEXT: ret;435 %val = call i32 @llvm.fshr.i32(i32 %a, i32 %b, i32 %c)436 ret i32 %val437}438 439define i32 @funnel_shift_left_32(i32 %a, i32 %b, i32 %c) {440; SM20-LABEL: funnel_shift_left_32(441; SM20: {442; SM20-NEXT: .reg .b32 %r<11>;443; SM20-EMPTY:444; SM20-NEXT: // %bb.0:445; SM20-NEXT: ld.param.b32 %r1, [funnel_shift_left_32_param_0];446; SM20-NEXT: ld.param.b32 %r2, [funnel_shift_left_32_param_2];447; SM20-NEXT: and.b32 %r3, %r2, 31;448; SM20-NEXT: shl.b32 %r4, %r1, %r3;449; SM20-NEXT: ld.param.b32 %r5, [funnel_shift_left_32_param_1];450; SM20-NEXT: shr.u32 %r6, %r5, 1;451; SM20-NEXT: not.b32 %r7, %r2;452; SM20-NEXT: and.b32 %r8, %r7, 31;453; SM20-NEXT: shr.u32 %r9, %r6, %r8;454; SM20-NEXT: or.b32 %r10, %r4, %r9;455; SM20-NEXT: st.param.b32 [func_retval0], %r10;456; SM20-NEXT: ret;457;458; SM35-LABEL: funnel_shift_left_32(459; SM35: {460; SM35-NEXT: .reg .b32 %r<5>;461; SM35-EMPTY:462; SM35-NEXT: // %bb.0:463; SM35-NEXT: ld.param.b32 %r1, [funnel_shift_left_32_param_0];464; SM35-NEXT: ld.param.b32 %r2, [funnel_shift_left_32_param_1];465; SM35-NEXT: ld.param.b32 %r3, [funnel_shift_left_32_param_2];466; SM35-NEXT: shf.l.wrap.b32 %r4, %r2, %r1, %r3;467; SM35-NEXT: st.param.b32 [func_retval0], %r4;468; SM35-NEXT: ret;469 %val = call i32 @llvm.fshl.i32(i32 %a, i32 %b, i32 %c)470 ret i32 %val471}472 473define i64 @funnel_shift_right_64(i64 %a, i64 %b, i64 %c) {474; SM20-LABEL: funnel_shift_right_64(475; SM20: {476; SM20-NEXT: .reg .b32 %r<5>;477; SM20-NEXT: .reg .b64 %rd<7>;478; SM20-EMPTY:479; SM20-NEXT: // %bb.0:480; SM20-NEXT: ld.param.b64 %rd1, [funnel_shift_right_64_param_0];481; SM20-NEXT: ld.param.b32 %r1, [funnel_shift_right_64_param_2];482; SM20-NEXT: and.b32 %r2, %r1, 63;483; SM20-NEXT: ld.param.b64 %rd2, [funnel_shift_right_64_param_1];484; SM20-NEXT: shr.u64 %rd3, %rd2, %r2;485; SM20-NEXT: shl.b64 %rd4, %rd1, 1;486; SM20-NEXT: not.b32 %r3, %r1;487; SM20-NEXT: and.b32 %r4, %r3, 63;488; SM20-NEXT: shl.b64 %rd5, %rd4, %r4;489; SM20-NEXT: or.b64 %rd6, %rd5, %rd3;490; SM20-NEXT: st.param.b64 [func_retval0], %rd6;491; SM20-NEXT: ret;492;493; SM35-LABEL: funnel_shift_right_64(494; SM35: {495; SM35-NEXT: .reg .b32 %r<5>;496; SM35-NEXT: .reg .b64 %rd<7>;497; SM35-EMPTY:498; SM35-NEXT: // %bb.0:499; SM35-NEXT: ld.param.b64 %rd1, [funnel_shift_right_64_param_0];500; SM35-NEXT: ld.param.b32 %r1, [funnel_shift_right_64_param_2];501; SM35-NEXT: and.b32 %r2, %r1, 63;502; SM35-NEXT: ld.param.b64 %rd2, [funnel_shift_right_64_param_1];503; SM35-NEXT: shr.u64 %rd3, %rd2, %r2;504; SM35-NEXT: shl.b64 %rd4, %rd1, 1;505; SM35-NEXT: not.b32 %r3, %r1;506; SM35-NEXT: and.b32 %r4, %r3, 63;507; SM35-NEXT: shl.b64 %rd5, %rd4, %r4;508; SM35-NEXT: or.b64 %rd6, %rd5, %rd3;509; SM35-NEXT: st.param.b64 [func_retval0], %rd6;510; SM35-NEXT: ret;511 %val = call i64 @llvm.fshr.i64(i64 %a, i64 %b, i64 %c)512 ret i64 %val513}514 515define i64 @funnel_shift_left_64(i64 %a, i64 %b, i64 %c) {516; SM20-LABEL: funnel_shift_left_64(517; SM20: {518; SM20-NEXT: .reg .b32 %r<5>;519; SM20-NEXT: .reg .b64 %rd<7>;520; SM20-EMPTY:521; SM20-NEXT: // %bb.0:522; SM20-NEXT: ld.param.b64 %rd1, [funnel_shift_left_64_param_0];523; SM20-NEXT: ld.param.b32 %r1, [funnel_shift_left_64_param_2];524; SM20-NEXT: and.b32 %r2, %r1, 63;525; SM20-NEXT: shl.b64 %rd2, %rd1, %r2;526; SM20-NEXT: ld.param.b64 %rd3, [funnel_shift_left_64_param_1];527; SM20-NEXT: shr.u64 %rd4, %rd3, 1;528; SM20-NEXT: not.b32 %r3, %r1;529; SM20-NEXT: and.b32 %r4, %r3, 63;530; SM20-NEXT: shr.u64 %rd5, %rd4, %r4;531; SM20-NEXT: or.b64 %rd6, %rd2, %rd5;532; SM20-NEXT: st.param.b64 [func_retval0], %rd6;533; SM20-NEXT: ret;534;535; SM35-LABEL: funnel_shift_left_64(536; SM35: {537; SM35-NEXT: .reg .b32 %r<5>;538; SM35-NEXT: .reg .b64 %rd<7>;539; SM35-EMPTY:540; SM35-NEXT: // %bb.0:541; SM35-NEXT: ld.param.b64 %rd1, [funnel_shift_left_64_param_0];542; SM35-NEXT: ld.param.b32 %r1, [funnel_shift_left_64_param_2];543; SM35-NEXT: and.b32 %r2, %r1, 63;544; SM35-NEXT: shl.b64 %rd2, %rd1, %r2;545; SM35-NEXT: ld.param.b64 %rd3, [funnel_shift_left_64_param_1];546; SM35-NEXT: shr.u64 %rd4, %rd3, 1;547; SM35-NEXT: not.b32 %r3, %r1;548; SM35-NEXT: and.b32 %r4, %r3, 63;549; SM35-NEXT: shr.u64 %rd5, %rd4, %r4;550; SM35-NEXT: or.b64 %rd6, %rd2, %rd5;551; SM35-NEXT: st.param.b64 [func_retval0], %rd6;552; SM35-NEXT: ret;553 %val = call i64 @llvm.fshl.i64(i64 %a, i64 %b, i64 %c)554 ret i64 %val555}556 557define i64 @fshl64_low_imm(i64 %a, i64 %b) {558; SM20-LABEL: fshl64_low_imm(559; SM20: {560; SM20-NEXT: .reg .b64 %rd<6>;561; SM20-EMPTY:562; SM20-NEXT: // %bb.0:563; SM20-NEXT: ld.param.b64 %rd1, [fshl64_low_imm_param_0];564; SM20-NEXT: ld.param.b64 %rd2, [fshl64_low_imm_param_1];565; SM20-NEXT: shr.u64 %rd3, %rd2, 59;566; SM20-NEXT: shl.b64 %rd4, %rd1, 5;567; SM20-NEXT: or.b64 %rd5, %rd4, %rd3;568; SM20-NEXT: st.param.b64 [func_retval0], %rd5;569; SM20-NEXT: ret;570;571; SM35-LABEL: fshl64_low_imm(572; SM35: {573; SM35-NEXT: .reg .b32 %r<7>;574; SM35-NEXT: .reg .b64 %rd<4>;575; SM35-EMPTY:576; SM35-NEXT: // %bb.0:577; SM35-NEXT: ld.param.b64 %rd1, [fshl64_low_imm_param_0];578; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;579; SM35-NEXT: ld.param.b64 %rd2, [fshl64_low_imm_param_1];580; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;581; SM35-NEXT: shf.l.wrap.b32 %r5, %r4, %r1, 5;582; SM35-NEXT: shf.l.wrap.b32 %r6, %r1, %r2, 5;583; SM35-NEXT: mov.b64 %rd3, {%r5, %r6};584; SM35-NEXT: st.param.b64 [func_retval0], %rd3;585; SM35-NEXT: ret;586 %val = call i64 @llvm.fshl.i64(i64 %a, i64 %b, i64 5)587 ret i64 %val588}589 590define i64 @fshl64_high_imm(i64 %a, i64 %b) {591; SM20-LABEL: fshl64_high_imm(592; SM20: {593; SM20-NEXT: .reg .b64 %rd<6>;594; SM20-EMPTY:595; SM20-NEXT: // %bb.0:596; SM20-NEXT: ld.param.b64 %rd1, [fshl64_high_imm_param_0];597; SM20-NEXT: ld.param.b64 %rd2, [fshl64_high_imm_param_1];598; SM20-NEXT: shr.u64 %rd3, %rd2, 9;599; SM20-NEXT: shl.b64 %rd4, %rd1, 55;600; SM20-NEXT: or.b64 %rd5, %rd4, %rd3;601; SM20-NEXT: st.param.b64 [func_retval0], %rd5;602; SM20-NEXT: ret;603;604; SM35-LABEL: fshl64_high_imm(605; SM35: {606; SM35-NEXT: .reg .b32 %r<7>;607; SM35-NEXT: .reg .b64 %rd<4>;608; SM35-EMPTY:609; SM35-NEXT: // %bb.0:610; SM35-NEXT: ld.param.b64 %rd1, [fshl64_high_imm_param_0];611; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;612; SM35-NEXT: ld.param.b64 %rd2, [fshl64_high_imm_param_1];613; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;614; SM35-NEXT: shf.l.wrap.b32 %r5, %r4, %r1, 23;615; SM35-NEXT: shf.l.wrap.b32 %r6, %r3, %r4, 23;616; SM35-NEXT: mov.b64 %rd3, {%r6, %r5};617; SM35-NEXT: st.param.b64 [func_retval0], %rd3;618; SM35-NEXT: ret;619 %val = call i64 @llvm.fshl.i64(i64 %a, i64 %b, i64 55)620 ret i64 %val621}622 623define i64 @fshl64_32_imm(i64 %a, i64 %b) {624; SM20-LABEL: fshl64_32_imm(625; SM20: {626; SM20-NEXT: .reg .b64 %rd<5>;627; SM20-EMPTY:628; SM20-NEXT: // %bb.0:629; SM20-NEXT: ld.param.b64 %rd1, [fshl64_32_imm_param_0];630; SM20-NEXT: shl.b64 %rd2, %rd1, 32;631; SM20-NEXT: ld.param.b32 %rd3, [fshl64_32_imm_param_1+4];632; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;633; SM20-NEXT: st.param.b64 [func_retval0], %rd4;634; SM20-NEXT: ret;635;636; SM35-LABEL: fshl64_32_imm(637; SM35: {638; SM35-NEXT: .reg .b32 %r<5>;639; SM35-NEXT: .reg .b64 %rd<4>;640; SM35-EMPTY:641; SM35-NEXT: // %bb.0:642; SM35-NEXT: ld.param.b64 %rd1, [fshl64_32_imm_param_0];643; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;644; SM35-NEXT: ld.param.b64 %rd2, [fshl64_32_imm_param_1];645; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;646; SM35-NEXT: mov.b64 %rd3, {%r4, %r1};647; SM35-NEXT: st.param.b64 [func_retval0], %rd3;648; SM35-NEXT: ret;649 %val = call i64 @llvm.fshl.i64(i64 %a, i64 %b, i64 32)650 ret i64 %val651}652 653define i64 @fshr64_low_imm(i64 %a, i64 %b) {654; SM20-LABEL: fshr64_low_imm(655; SM20: {656; SM20-NEXT: .reg .b64 %rd<6>;657; SM20-EMPTY:658; SM20-NEXT: // %bb.0:659; SM20-NEXT: ld.param.b64 %rd1, [fshr64_low_imm_param_0];660; SM20-NEXT: ld.param.b64 %rd2, [fshr64_low_imm_param_1];661; SM20-NEXT: shr.u64 %rd3, %rd2, 31;662; SM20-NEXT: shl.b64 %rd4, %rd1, 33;663; SM20-NEXT: or.b64 %rd5, %rd4, %rd3;664; SM20-NEXT: st.param.b64 [func_retval0], %rd5;665; SM20-NEXT: ret;666;667; SM35-LABEL: fshr64_low_imm(668; SM35: {669; SM35-NEXT: .reg .b32 %r<7>;670; SM35-NEXT: .reg .b64 %rd<4>;671; SM35-EMPTY:672; SM35-NEXT: // %bb.0:673; SM35-NEXT: ld.param.b64 %rd1, [fshr64_low_imm_param_0];674; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;675; SM35-NEXT: ld.param.b64 %rd2, [fshr64_low_imm_param_1];676; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;677; SM35-NEXT: shf.r.wrap.b32 %r5, %r4, %r1, 31;678; SM35-NEXT: shf.r.wrap.b32 %r6, %r3, %r4, 31;679; SM35-NEXT: mov.b64 %rd3, {%r6, %r5};680; SM35-NEXT: st.param.b64 [func_retval0], %rd3;681; SM35-NEXT: ret;682 %val = call i64 @llvm.fshr.i64(i64 %a, i64 %b, i64 31)683 ret i64 %val684}685 686define i64 @fshr64_high_imm(i64 %a, i64 %b) {687; SM20-LABEL: fshr64_high_imm(688; SM20: {689; SM20-NEXT: .reg .b64 %rd<6>;690; SM20-EMPTY:691; SM20-NEXT: // %bb.0:692; SM20-NEXT: ld.param.b64 %rd1, [fshr64_high_imm_param_0];693; SM20-NEXT: ld.param.b64 %rd2, [fshr64_high_imm_param_1];694; SM20-NEXT: shr.u64 %rd3, %rd2, 33;695; SM20-NEXT: shl.b64 %rd4, %rd1, 31;696; SM20-NEXT: or.b64 %rd5, %rd4, %rd3;697; SM20-NEXT: st.param.b64 [func_retval0], %rd5;698; SM20-NEXT: ret;699;700; SM35-LABEL: fshr64_high_imm(701; SM35: {702; SM35-NEXT: .reg .b32 %r<7>;703; SM35-NEXT: .reg .b64 %rd<4>;704; SM35-EMPTY:705; SM35-NEXT: // %bb.0:706; SM35-NEXT: ld.param.b64 %rd1, [fshr64_high_imm_param_0];707; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;708; SM35-NEXT: ld.param.b64 %rd2, [fshr64_high_imm_param_1];709; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;710; SM35-NEXT: shf.r.wrap.b32 %r5, %r4, %r1, 1;711; SM35-NEXT: shf.r.wrap.b32 %r6, %r1, %r2, 1;712; SM35-NEXT: mov.b64 %rd3, {%r5, %r6};713; SM35-NEXT: st.param.b64 [func_retval0], %rd3;714; SM35-NEXT: ret;715 %val = call i64 @llvm.fshr.i64(i64 %a, i64 %b, i64 33)716 ret i64 %val717}718 719define i64 @fshr64_32_imm(i64 %a, i64 %b) {720; SM20-LABEL: fshr64_32_imm(721; SM20: {722; SM20-NEXT: .reg .b64 %rd<5>;723; SM20-EMPTY:724; SM20-NEXT: // %bb.0:725; SM20-NEXT: ld.param.b64 %rd1, [fshr64_32_imm_param_0];726; SM20-NEXT: shl.b64 %rd2, %rd1, 32;727; SM20-NEXT: ld.param.b32 %rd3, [fshr64_32_imm_param_1+4];728; SM20-NEXT: or.b64 %rd4, %rd2, %rd3;729; SM20-NEXT: st.param.b64 [func_retval0], %rd4;730; SM20-NEXT: ret;731;732; SM35-LABEL: fshr64_32_imm(733; SM35: {734; SM35-NEXT: .reg .b32 %r<5>;735; SM35-NEXT: .reg .b64 %rd<4>;736; SM35-EMPTY:737; SM35-NEXT: // %bb.0:738; SM35-NEXT: ld.param.b64 %rd1, [fshr64_32_imm_param_0];739; SM35-NEXT: mov.b64 {%r1, %r2}, %rd1;740; SM35-NEXT: ld.param.b64 %rd2, [fshr64_32_imm_param_1];741; SM35-NEXT: mov.b64 {%r3, %r4}, %rd2;742; SM35-NEXT: mov.b64 %rd3, {%r4, %r1};743; SM35-NEXT: st.param.b64 [func_retval0], %rd3;744; SM35-NEXT: ret;745 %val = call i64 @llvm.fshr.i64(i64 %a, i64 %b, i64 32)746 ret i64 %val747}748