2571 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_80 -mattr=+ptx70 -O0 \3; RUN: -disable-post-ra -verify-machineinstrs \4; RUN: | FileCheck -check-prefixes CHECK,CHECK-SM80 %s5; RUN: %if ptxas-sm_80 && ptxas-isa-7.0 %{ llc < %s -mcpu=sm_80 -mattr=+ptx70 -O0 \6; RUN: -disable-post-ra -verify-machineinstrs \7; RUN: | %ptxas-verify -arch=sm_80 %}8; RUN: llc < %s -mcpu=sm_100 -mattr=+ptx88 -O0 \9; RUN: -disable-post-ra -verify-machineinstrs \10; RUN: | FileCheck -check-prefixes CHECK,CHECK-SM100 %s11; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -mcpu=sm_100 -mattr=+ptx88 -O0 \12; RUN: -disable-post-ra -verify-machineinstrs \13; RUN: | %ptxas-verify -arch=sm_100 %}14target triple = "nvptx64-nvidia-cuda"15target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"16 17define half @reduce_fadd_half(<8 x half> %in) {18; CHECK-LABEL: reduce_fadd_half(19; CHECK: {20; CHECK-NEXT: .reg .b16 %rs<18>;21; CHECK-NEXT: .reg .b32 %r<5>;22; CHECK-EMPTY:23; CHECK-NEXT: // %bb.0:24; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_half_param_0];25; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;26; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;27; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r2;28; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r1;29; CHECK-NEXT: mov.b16 %rs9, 0x0000;30; CHECK-NEXT: add.rn.f16 %rs10, %rs7, %rs9;31; CHECK-NEXT: add.rn.f16 %rs11, %rs10, %rs8;32; CHECK-NEXT: add.rn.f16 %rs12, %rs11, %rs5;33; CHECK-NEXT: add.rn.f16 %rs13, %rs12, %rs6;34; CHECK-NEXT: add.rn.f16 %rs14, %rs13, %rs3;35; CHECK-NEXT: add.rn.f16 %rs15, %rs14, %rs4;36; CHECK-NEXT: add.rn.f16 %rs16, %rs15, %rs1;37; CHECK-NEXT: add.rn.f16 %rs17, %rs16, %rs2;38; CHECK-NEXT: st.param.b16 [func_retval0], %rs17;39; CHECK-NEXT: ret;40 %res = call half @llvm.vector.reduce.fadd(half 0.0, <8 x half> %in)41 ret half %res42}43 44define half @reduce_fadd_half_reassoc(<8 x half> %in) {45; CHECK-LABEL: reduce_fadd_half_reassoc(46; CHECK: {47; CHECK-NEXT: .reg .b16 %rs<6>;48; CHECK-NEXT: .reg .b32 %r<8>;49; CHECK-EMPTY:50; CHECK-NEXT: // %bb.0:51; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_half_reassoc_param_0];52; CHECK-NEXT: add.rn.f16x2 %r5, %r2, %r4;53; CHECK-NEXT: add.rn.f16x2 %r6, %r1, %r3;54; CHECK-NEXT: add.rn.f16x2 %r7, %r6, %r5;55; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;56; CHECK-NEXT: add.rn.f16 %rs3, %rs1, %rs2;57; CHECK-NEXT: mov.b16 %rs4, 0x0000;58; CHECK-NEXT: add.rn.f16 %rs5, %rs3, %rs4;59; CHECK-NEXT: st.param.b16 [func_retval0], %rs5;60; CHECK-NEXT: ret;61 %res = call reassoc half @llvm.vector.reduce.fadd(half 0.0, <8 x half> %in)62 ret half %res63}64 65define half @reduce_fadd_half_reassoc_nonpow2(<7 x half> %in) {66; CHECK-LABEL: reduce_fadd_half_reassoc_nonpow2(67; CHECK: {68; CHECK-NEXT: .reg .b16 %rs<16>;69; CHECK-EMPTY:70; CHECK-NEXT: // %bb.0:71; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fadd_half_reassoc_nonpow2_param_0+12];72; CHECK-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_fadd_half_reassoc_nonpow2_param_0+8];73; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_fadd_half_reassoc_nonpow2_param_0];74; CHECK-NEXT: mov.b16 %rs8, 0x0000;75; CHECK-NEXT: add.rn.f16 %rs9, %rs1, %rs8;76; CHECK-NEXT: add.rn.f16 %rs10, %rs9, %rs2;77; CHECK-NEXT: add.rn.f16 %rs11, %rs10, %rs3;78; CHECK-NEXT: add.rn.f16 %rs12, %rs11, %rs4;79; CHECK-NEXT: add.rn.f16 %rs13, %rs12, %rs5;80; CHECK-NEXT: add.rn.f16 %rs14, %rs13, %rs6;81; CHECK-NEXT: add.rn.f16 %rs15, %rs14, %rs7;82; CHECK-NEXT: st.param.b16 [func_retval0], %rs15;83; CHECK-NEXT: ret;84 %res = call half @llvm.vector.reduce.fadd(half 0.0, <7 x half> %in)85 ret half %res86}87 88define float @reduce_fadd_float(<8 x float> %in) {89; CHECK-SM80-LABEL: reduce_fadd_float(90; CHECK-SM80: {91; CHECK-SM80-NEXT: .reg .b32 %r<17>;92; CHECK-SM80-EMPTY:93; CHECK-SM80-NEXT: // %bb.0:94; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_param_0+16];95; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_param_0];96; CHECK-SM80-NEXT: add.rn.f32 %r9, %r1, 0f00000000;97; CHECK-SM80-NEXT: add.rn.f32 %r10, %r9, %r2;98; CHECK-SM80-NEXT: add.rn.f32 %r11, %r10, %r3;99; CHECK-SM80-NEXT: add.rn.f32 %r12, %r11, %r4;100; CHECK-SM80-NEXT: add.rn.f32 %r13, %r12, %r5;101; CHECK-SM80-NEXT: add.rn.f32 %r14, %r13, %r6;102; CHECK-SM80-NEXT: add.rn.f32 %r15, %r14, %r7;103; CHECK-SM80-NEXT: add.rn.f32 %r16, %r15, %r8;104; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r16;105; CHECK-SM80-NEXT: ret;106;107; CHECK-SM100-LABEL: reduce_fadd_float(108; CHECK-SM100: {109; CHECK-SM100-NEXT: .reg .b32 %r<17>;110; CHECK-SM100-NEXT: .reg .b64 %rd<5>;111; CHECK-SM100-EMPTY:112; CHECK-SM100-NEXT: // %bb.0:113; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fadd_float_param_0+16];114; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;115; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd3;116; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fadd_float_param_0];117; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd2;118; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd1;119; CHECK-SM100-NEXT: add.rn.f32 %r9, %r7, 0f00000000;120; CHECK-SM100-NEXT: add.rn.f32 %r10, %r9, %r8;121; CHECK-SM100-NEXT: add.rn.f32 %r11, %r10, %r5;122; CHECK-SM100-NEXT: add.rn.f32 %r12, %r11, %r6;123; CHECK-SM100-NEXT: add.rn.f32 %r13, %r12, %r3;124; CHECK-SM100-NEXT: add.rn.f32 %r14, %r13, %r4;125; CHECK-SM100-NEXT: add.rn.f32 %r15, %r14, %r1;126; CHECK-SM100-NEXT: add.rn.f32 %r16, %r15, %r2;127; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r16;128; CHECK-SM100-NEXT: ret;129 %res = call float @llvm.vector.reduce.fadd(float 0.0, <8 x float> %in)130 ret float %res131}132 133define float @reduce_fadd_float_reassoc(<8 x float> %in) {134; CHECK-SM80-LABEL: reduce_fadd_float_reassoc(135; CHECK-SM80: {136; CHECK-SM80-NEXT: .reg .b32 %r<17>;137; CHECK-SM80-EMPTY:138; CHECK-SM80-NEXT: // %bb.0:139; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_reassoc_param_0+16];140; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_reassoc_param_0];141; CHECK-SM80-NEXT: add.rn.f32 %r9, %r4, %r8;142; CHECK-SM80-NEXT: add.rn.f32 %r10, %r2, %r6;143; CHECK-SM80-NEXT: add.rn.f32 %r11, %r10, %r9;144; CHECK-SM80-NEXT: add.rn.f32 %r12, %r3, %r7;145; CHECK-SM80-NEXT: add.rn.f32 %r13, %r1, %r5;146; CHECK-SM80-NEXT: add.rn.f32 %r14, %r13, %r12;147; CHECK-SM80-NEXT: add.rn.f32 %r15, %r14, %r11;148; CHECK-SM80-NEXT: add.rn.f32 %r16, %r15, 0f00000000;149; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r16;150; CHECK-SM80-NEXT: ret;151;152; CHECK-SM100-LABEL: reduce_fadd_float_reassoc(153; CHECK-SM100: {154; CHECK-SM100-NEXT: .reg .b32 %r<5>;155; CHECK-SM100-NEXT: .reg .b64 %rd<8>;156; CHECK-SM100-EMPTY:157; CHECK-SM100-NEXT: // %bb.0:158; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fadd_float_reassoc_param_0+16];159; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fadd_float_reassoc_param_0];160; CHECK-SM100-NEXT: add.rn.f32x2 %rd5, %rd2, %rd4;161; CHECK-SM100-NEXT: add.rn.f32x2 %rd6, %rd1, %rd3;162; CHECK-SM100-NEXT: add.rn.f32x2 %rd7, %rd6, %rd5;163; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd7;164; CHECK-SM100-NEXT: add.rn.f32 %r3, %r1, %r2;165; CHECK-SM100-NEXT: add.rn.f32 %r4, %r3, 0f00000000;166; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r4;167; CHECK-SM100-NEXT: ret;168 %res = call reassoc float @llvm.vector.reduce.fadd(float 0.0, <8 x float> %in)169 ret float %res170}171 172define float @reduce_fadd_float_reassoc_nonpow2(<7 x float> %in) {173; CHECK-SM80-LABEL: reduce_fadd_float_reassoc_nonpow2(174; CHECK-SM80: {175; CHECK-SM80-NEXT: .reg .b32 %r<15>;176; CHECK-SM80-EMPTY:177; CHECK-SM80-NEXT: // %bb.0:178; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fadd_float_reassoc_nonpow2_param_0+24];179; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fadd_float_reassoc_nonpow2_param_0+16];180; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_reassoc_nonpow2_param_0];181; CHECK-SM80-NEXT: add.rn.f32 %r8, %r3, %r7;182; CHECK-SM80-NEXT: add.rn.f32 %r9, %r1, %r5;183; CHECK-SM80-NEXT: add.rn.f32 %r10, %r9, %r8;184; CHECK-SM80-NEXT: add.rn.f32 %r11, %r2, %r6;185; CHECK-SM80-NEXT: add.rn.f32 %r12, %r11, %r4;186; CHECK-SM80-NEXT: add.rn.f32 %r13, %r10, %r12;187; CHECK-SM80-NEXT: add.rn.f32 %r14, %r13, 0f00000000;188; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r14;189; CHECK-SM80-NEXT: ret;190;191; CHECK-SM100-LABEL: reduce_fadd_float_reassoc_nonpow2(192; CHECK-SM100: {193; CHECK-SM100-NEXT: .reg .b32 %r<13>;194; CHECK-SM100-NEXT: .reg .b64 %rd<8>;195; CHECK-SM100-EMPTY:196; CHECK-SM100-NEXT: // %bb.0:197; CHECK-SM100-NEXT: ld.param.b64 %rd1, [reduce_fadd_float_reassoc_nonpow2_param_0+16];198; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd1;199; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [reduce_fadd_float_reassoc_nonpow2_param_0];200; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd3;201; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd2;202; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fadd_float_reassoc_nonpow2_param_0+24];203; CHECK-SM100-NEXT: add.rn.f32x2 %rd4, %rd2, %rd1;204; CHECK-SM100-NEXT: mov.b32 %r8, 0f80000000;205; CHECK-SM100-NEXT: mov.b64 %rd5, {%r7, %r8};206; CHECK-SM100-NEXT: add.rn.f32x2 %rd6, %rd3, %rd5;207; CHECK-SM100-NEXT: add.rn.f32x2 %rd7, %rd4, %rd6;208; CHECK-SM100-NEXT: mov.b64 {%r9, %r10}, %rd7;209; CHECK-SM100-NEXT: add.rn.f32 %r11, %r9, %r10;210; CHECK-SM100-NEXT: add.rn.f32 %r12, %r11, 0f00000000;211; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;212; CHECK-SM100-NEXT: ret;213 %res = call reassoc float @llvm.vector.reduce.fadd(float 0.0, <7 x float> %in)214 ret float %res215}216 217define half @reduce_fmul_half(<8 x half> %in) {218; CHECK-LABEL: reduce_fmul_half(219; CHECK: {220; CHECK-NEXT: .reg .b16 %rs<16>;221; CHECK-NEXT: .reg .b32 %r<5>;222; CHECK-EMPTY:223; CHECK-NEXT: // %bb.0:224; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_half_param_0];225; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;226; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;227; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r2;228; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r1;229; CHECK-NEXT: mul.rn.f16 %rs9, %rs7, %rs8;230; CHECK-NEXT: mul.rn.f16 %rs10, %rs9, %rs5;231; CHECK-NEXT: mul.rn.f16 %rs11, %rs10, %rs6;232; CHECK-NEXT: mul.rn.f16 %rs12, %rs11, %rs3;233; CHECK-NEXT: mul.rn.f16 %rs13, %rs12, %rs4;234; CHECK-NEXT: mul.rn.f16 %rs14, %rs13, %rs1;235; CHECK-NEXT: mul.rn.f16 %rs15, %rs14, %rs2;236; CHECK-NEXT: st.param.b16 [func_retval0], %rs15;237; CHECK-NEXT: ret;238 %res = call half @llvm.vector.reduce.fmul(half 1.0, <8 x half> %in)239 ret half %res240}241 242define half @reduce_fmul_half_reassoc(<8 x half> %in) {243; CHECK-LABEL: reduce_fmul_half_reassoc(244; CHECK: {245; CHECK-NEXT: .reg .b16 %rs<4>;246; CHECK-NEXT: .reg .b32 %r<8>;247; CHECK-EMPTY:248; CHECK-NEXT: // %bb.0:249; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_half_reassoc_param_0];250; CHECK-NEXT: mul.rn.f16x2 %r5, %r2, %r4;251; CHECK-NEXT: mul.rn.f16x2 %r6, %r1, %r3;252; CHECK-NEXT: mul.rn.f16x2 %r7, %r6, %r5;253; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;254; CHECK-NEXT: mul.rn.f16 %rs3, %rs1, %rs2;255; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;256; CHECK-NEXT: ret;257 %res = call reassoc half @llvm.vector.reduce.fmul(half 1.0, <8 x half> %in)258 ret half %res259}260 261define half @reduce_fmul_half_reassoc_nonpow2(<7 x half> %in) {262; CHECK-LABEL: reduce_fmul_half_reassoc_nonpow2(263; CHECK: {264; CHECK-NEXT: .reg .b16 %rs<12>;265; CHECK-NEXT: .reg .b32 %r<8>;266; CHECK-EMPTY:267; CHECK-NEXT: // %bb.0:268; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmul_half_reassoc_nonpow2_param_0+8];269; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;270; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmul_half_reassoc_nonpow2_param_0];271; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;272; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;273; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmul_half_reassoc_nonpow2_param_0+12];274; CHECK-NEXT: mul.rn.f16x2 %r4, %r2, %r1;275; CHECK-NEXT: mov.b16 %rs8, 0x3C00;276; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};277; CHECK-NEXT: mul.rn.f16x2 %r6, %r3, %r5;278; CHECK-NEXT: mul.rn.f16x2 %r7, %r4, %r6;279; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;280; CHECK-NEXT: mul.rn.f16 %rs11, %rs9, %rs10;281; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;282; CHECK-NEXT: ret;283 %res = call reassoc half @llvm.vector.reduce.fmul(half 1.0, <7 x half> %in)284 ret half %res285}286 287define float @reduce_fmul_float(<8 x float> %in) {288; CHECK-SM80-LABEL: reduce_fmul_float(289; CHECK-SM80: {290; CHECK-SM80-NEXT: .reg .b32 %r<16>;291; CHECK-SM80-EMPTY:292; CHECK-SM80-NEXT: // %bb.0:293; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_param_0+16];294; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_param_0];295; CHECK-SM80-NEXT: mul.rn.f32 %r9, %r1, %r2;296; CHECK-SM80-NEXT: mul.rn.f32 %r10, %r9, %r3;297; CHECK-SM80-NEXT: mul.rn.f32 %r11, %r10, %r4;298; CHECK-SM80-NEXT: mul.rn.f32 %r12, %r11, %r5;299; CHECK-SM80-NEXT: mul.rn.f32 %r13, %r12, %r6;300; CHECK-SM80-NEXT: mul.rn.f32 %r14, %r13, %r7;301; CHECK-SM80-NEXT: mul.rn.f32 %r15, %r14, %r8;302; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;303; CHECK-SM80-NEXT: ret;304;305; CHECK-SM100-LABEL: reduce_fmul_float(306; CHECK-SM100: {307; CHECK-SM100-NEXT: .reg .b32 %r<16>;308; CHECK-SM100-NEXT: .reg .b64 %rd<5>;309; CHECK-SM100-EMPTY:310; CHECK-SM100-NEXT: // %bb.0:311; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmul_float_param_0+16];312; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;313; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd3;314; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmul_float_param_0];315; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd2;316; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd1;317; CHECK-SM100-NEXT: mul.rn.f32 %r9, %r7, %r8;318; CHECK-SM100-NEXT: mul.rn.f32 %r10, %r9, %r5;319; CHECK-SM100-NEXT: mul.rn.f32 %r11, %r10, %r6;320; CHECK-SM100-NEXT: mul.rn.f32 %r12, %r11, %r3;321; CHECK-SM100-NEXT: mul.rn.f32 %r13, %r12, %r4;322; CHECK-SM100-NEXT: mul.rn.f32 %r14, %r13, %r1;323; CHECK-SM100-NEXT: mul.rn.f32 %r15, %r14, %r2;324; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;325; CHECK-SM100-NEXT: ret;326 %res = call float @llvm.vector.reduce.fmul(float 1.0, <8 x float> %in)327 ret float %res328}329 330define float @reduce_fmul_float_reassoc(<8 x float> %in) {331; CHECK-SM80-LABEL: reduce_fmul_float_reassoc(332; CHECK-SM80: {333; CHECK-SM80-NEXT: .reg .b32 %r<16>;334; CHECK-SM80-EMPTY:335; CHECK-SM80-NEXT: // %bb.0:336; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_reassoc_param_0+16];337; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_reassoc_param_0];338; CHECK-SM80-NEXT: mul.rn.f32 %r9, %r4, %r8;339; CHECK-SM80-NEXT: mul.rn.f32 %r10, %r2, %r6;340; CHECK-SM80-NEXT: mul.rn.f32 %r11, %r10, %r9;341; CHECK-SM80-NEXT: mul.rn.f32 %r12, %r3, %r7;342; CHECK-SM80-NEXT: mul.rn.f32 %r13, %r1, %r5;343; CHECK-SM80-NEXT: mul.rn.f32 %r14, %r13, %r12;344; CHECK-SM80-NEXT: mul.rn.f32 %r15, %r14, %r11;345; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;346; CHECK-SM80-NEXT: ret;347;348; CHECK-SM100-LABEL: reduce_fmul_float_reassoc(349; CHECK-SM100: {350; CHECK-SM100-NEXT: .reg .b32 %r<4>;351; CHECK-SM100-NEXT: .reg .b64 %rd<8>;352; CHECK-SM100-EMPTY:353; CHECK-SM100-NEXT: // %bb.0:354; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmul_float_reassoc_param_0+16];355; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmul_float_reassoc_param_0];356; CHECK-SM100-NEXT: mul.rn.f32x2 %rd5, %rd2, %rd4;357; CHECK-SM100-NEXT: mul.rn.f32x2 %rd6, %rd1, %rd3;358; CHECK-SM100-NEXT: mul.rn.f32x2 %rd7, %rd6, %rd5;359; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd7;360; CHECK-SM100-NEXT: mul.rn.f32 %r3, %r1, %r2;361; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r3;362; CHECK-SM100-NEXT: ret;363 %res = call reassoc float @llvm.vector.reduce.fmul(float 1.0, <8 x float> %in)364 ret float %res365}366 367define float @reduce_fmul_float_reassoc_nonpow2(<7 x float> %in) {368; CHECK-SM80-LABEL: reduce_fmul_float_reassoc_nonpow2(369; CHECK-SM80: {370; CHECK-SM80-NEXT: .reg .b32 %r<14>;371; CHECK-SM80-EMPTY:372; CHECK-SM80-NEXT: // %bb.0:373; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmul_float_reassoc_nonpow2_param_0+24];374; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmul_float_reassoc_nonpow2_param_0+16];375; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_reassoc_nonpow2_param_0];376; CHECK-SM80-NEXT: mul.rn.f32 %r8, %r3, %r7;377; CHECK-SM80-NEXT: mul.rn.f32 %r9, %r1, %r5;378; CHECK-SM80-NEXT: mul.rn.f32 %r10, %r9, %r8;379; CHECK-SM80-NEXT: mul.rn.f32 %r11, %r2, %r6;380; CHECK-SM80-NEXT: mul.rn.f32 %r12, %r11, %r4;381; CHECK-SM80-NEXT: mul.rn.f32 %r13, %r10, %r12;382; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;383; CHECK-SM80-NEXT: ret;384;385; CHECK-SM100-LABEL: reduce_fmul_float_reassoc_nonpow2(386; CHECK-SM100: {387; CHECK-SM100-NEXT: .reg .b32 %r<12>;388; CHECK-SM100-NEXT: .reg .b64 %rd<8>;389; CHECK-SM100-EMPTY:390; CHECK-SM100-NEXT: // %bb.0:391; CHECK-SM100-NEXT: ld.param.b64 %rd1, [reduce_fmul_float_reassoc_nonpow2_param_0+16];392; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd1;393; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [reduce_fmul_float_reassoc_nonpow2_param_0];394; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd3;395; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd2;396; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmul_float_reassoc_nonpow2_param_0+24];397; CHECK-SM100-NEXT: mul.rn.f32x2 %rd4, %rd2, %rd1;398; CHECK-SM100-NEXT: mov.b32 %r8, 0f3F800000;399; CHECK-SM100-NEXT: mov.b64 %rd5, {%r7, %r8};400; CHECK-SM100-NEXT: mul.rn.f32x2 %rd6, %rd3, %rd5;401; CHECK-SM100-NEXT: mul.rn.f32x2 %rd7, %rd4, %rd6;402; CHECK-SM100-NEXT: mov.b64 {%r9, %r10}, %rd7;403; CHECK-SM100-NEXT: mul.rn.f32 %r11, %r9, %r10;404; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r11;405; CHECK-SM100-NEXT: ret;406 %res = call reassoc float @llvm.vector.reduce.fmul(float 1.0, <7 x float> %in)407 ret float %res408}409 410define half @reduce_fmax_half(<8 x half> %in) {411; CHECK-LABEL: reduce_fmax_half(412; CHECK: {413; CHECK-NEXT: .reg .b16 %rs<4>;414; CHECK-NEXT: .reg .b32 %r<8>;415; CHECK-EMPTY:416; CHECK-NEXT: // %bb.0:417; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_half_param_0];418; CHECK-NEXT: max.f16x2 %r5, %r2, %r4;419; CHECK-NEXT: max.f16x2 %r6, %r1, %r3;420; CHECK-NEXT: max.f16x2 %r7, %r6, %r5;421; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;422; CHECK-NEXT: max.f16 %rs3, %rs1, %rs2;423; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;424; CHECK-NEXT: ret;425 %res = call half @llvm.vector.reduce.fmax(<8 x half> %in)426 ret half %res427}428 429define half @reduce_fmax_half_reassoc(<8 x half> %in) {430; CHECK-LABEL: reduce_fmax_half_reassoc(431; CHECK: {432; CHECK-NEXT: .reg .b16 %rs<4>;433; CHECK-NEXT: .reg .b32 %r<8>;434; CHECK-EMPTY:435; CHECK-NEXT: // %bb.0:436; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_half_reassoc_param_0];437; CHECK-NEXT: max.f16x2 %r5, %r2, %r4;438; CHECK-NEXT: max.f16x2 %r6, %r1, %r3;439; CHECK-NEXT: max.f16x2 %r7, %r6, %r5;440; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;441; CHECK-NEXT: max.f16 %rs3, %rs1, %rs2;442; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;443; CHECK-NEXT: ret;444 %res = call reassoc half @llvm.vector.reduce.fmax(<8 x half> %in)445 ret half %res446}447 448define half @reduce_fmax_half_reassoc_nonpow2(<7 x half> %in) {449; CHECK-LABEL: reduce_fmax_half_reassoc_nonpow2(450; CHECK: {451; CHECK-NEXT: .reg .b16 %rs<12>;452; CHECK-NEXT: .reg .b32 %r<8>;453; CHECK-EMPTY:454; CHECK-NEXT: // %bb.0:455; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmax_half_reassoc_nonpow2_param_0+8];456; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;457; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmax_half_reassoc_nonpow2_param_0];458; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;459; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;460; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmax_half_reassoc_nonpow2_param_0+12];461; CHECK-NEXT: max.f16x2 %r4, %r2, %r1;462; CHECK-NEXT: mov.b16 %rs8, 0xFE00;463; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};464; CHECK-NEXT: max.f16x2 %r6, %r3, %r5;465; CHECK-NEXT: max.f16x2 %r7, %r4, %r6;466; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;467; CHECK-NEXT: max.f16 %rs11, %rs9, %rs10;468; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;469; CHECK-NEXT: ret;470 %res = call reassoc half @llvm.vector.reduce.fmax(<7 x half> %in)471 ret half %res472}473 474define half @reduce_fmax_half_nnan(<8 x half> %in) {475; CHECK-LABEL: reduce_fmax_half_nnan(476; CHECK: {477; CHECK-NEXT: .reg .b16 %rs<4>;478; CHECK-NEXT: .reg .b32 %r<8>;479; CHECK-EMPTY:480; CHECK-NEXT: // %bb.0:481; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_half_nnan_param_0];482; CHECK-NEXT: max.f16x2 %r5, %r2, %r4;483; CHECK-NEXT: max.f16x2 %r6, %r1, %r3;484; CHECK-NEXT: max.f16x2 %r7, %r6, %r5;485; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;486; CHECK-NEXT: max.f16 %rs3, %rs1, %rs2;487; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;488; CHECK-NEXT: ret;489 %res = call nnan half @llvm.vector.reduce.fmax(<8 x half> %in)490 ret half %res491}492 493define half @reduce_fmax_half_nnan_nonpow2(<7 x half> %in) {494; CHECK-LABEL: reduce_fmax_half_nnan_nonpow2(495; CHECK: {496; CHECK-NEXT: .reg .b16 %rs<12>;497; CHECK-NEXT: .reg .b32 %r<8>;498; CHECK-EMPTY:499; CHECK-NEXT: // %bb.0:500; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmax_half_nnan_nonpow2_param_0+8];501; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;502; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmax_half_nnan_nonpow2_param_0];503; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;504; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;505; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmax_half_nnan_nonpow2_param_0+12];506; CHECK-NEXT: max.f16x2 %r4, %r2, %r1;507; CHECK-NEXT: mov.b16 %rs8, 0xFC00;508; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};509; CHECK-NEXT: max.f16x2 %r6, %r3, %r5;510; CHECK-NEXT: max.f16x2 %r7, %r4, %r6;511; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;512; CHECK-NEXT: max.f16 %rs11, %rs9, %rs10;513; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;514; CHECK-NEXT: ret;515 %res = call nnan half @llvm.vector.reduce.fmax(<7 x half> %in)516 ret half %res517}518 519define float @reduce_fmax_float(<8 x float> %in) {520; CHECK-SM80-LABEL: reduce_fmax_float(521; CHECK-SM80: {522; CHECK-SM80-NEXT: .reg .b32 %r<16>;523; CHECK-SM80-EMPTY:524; CHECK-SM80-NEXT: // %bb.0:525; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_param_0+16];526; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_param_0];527; CHECK-SM80-NEXT: max.f32 %r9, %r7, %r8;528; CHECK-SM80-NEXT: max.f32 %r10, %r5, %r6;529; CHECK-SM80-NEXT: max.f32 %r11, %r10, %r9;530; CHECK-SM80-NEXT: max.f32 %r12, %r3, %r4;531; CHECK-SM80-NEXT: max.f32 %r13, %r1, %r2;532; CHECK-SM80-NEXT: max.f32 %r14, %r13, %r12;533; CHECK-SM80-NEXT: max.f32 %r15, %r14, %r11;534; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;535; CHECK-SM80-NEXT: ret;536;537; CHECK-SM100-LABEL: reduce_fmax_float(538; CHECK-SM100: {539; CHECK-SM100-NEXT: .reg .b32 %r<13>;540; CHECK-SM100-NEXT: .reg .b64 %rd<5>;541; CHECK-SM100-EMPTY:542; CHECK-SM100-NEXT: // %bb.0:543; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmax_float_param_0+16];544; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;545; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmax_float_param_0];546; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;547; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;548; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;549; CHECK-SM100-NEXT: max.f32 %r9, %r8, %r5, %r6;550; CHECK-SM100-NEXT: max.f32 %r10, %r3, %r4, %r7;551; CHECK-SM100-NEXT: max.f32 %r11, %r10, %r9, %r1;552; CHECK-SM100-NEXT: max.f32 %r12, %r11, %r2;553; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;554; CHECK-SM100-NEXT: ret;555 %res = call float @llvm.vector.reduce.fmax(<8 x float> %in)556 ret float %res557}558 559define float @reduce_fmax_float_reassoc(<8 x float> %in) {560; CHECK-SM80-LABEL: reduce_fmax_float_reassoc(561; CHECK-SM80: {562; CHECK-SM80-NEXT: .reg .b32 %r<16>;563; CHECK-SM80-EMPTY:564; CHECK-SM80-NEXT: // %bb.0:565; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_reassoc_param_0+16];566; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_reassoc_param_0];567; CHECK-SM80-NEXT: max.f32 %r9, %r7, %r8;568; CHECK-SM80-NEXT: max.f32 %r10, %r5, %r6;569; CHECK-SM80-NEXT: max.f32 %r11, %r10, %r9;570; CHECK-SM80-NEXT: max.f32 %r12, %r3, %r4;571; CHECK-SM80-NEXT: max.f32 %r13, %r1, %r2;572; CHECK-SM80-NEXT: max.f32 %r14, %r13, %r12;573; CHECK-SM80-NEXT: max.f32 %r15, %r14, %r11;574; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;575; CHECK-SM80-NEXT: ret;576;577; CHECK-SM100-LABEL: reduce_fmax_float_reassoc(578; CHECK-SM100: {579; CHECK-SM100-NEXT: .reg .b32 %r<13>;580; CHECK-SM100-NEXT: .reg .b64 %rd<5>;581; CHECK-SM100-EMPTY:582; CHECK-SM100-NEXT: // %bb.0:583; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmax_float_reassoc_param_0+16];584; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;585; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmax_float_reassoc_param_0];586; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;587; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;588; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;589; CHECK-SM100-NEXT: max.f32 %r9, %r8, %r5, %r6;590; CHECK-SM100-NEXT: max.f32 %r10, %r3, %r4, %r7;591; CHECK-SM100-NEXT: max.f32 %r11, %r10, %r9, %r1;592; CHECK-SM100-NEXT: max.f32 %r12, %r11, %r2;593; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;594; CHECK-SM100-NEXT: ret;595 %res = call reassoc float @llvm.vector.reduce.fmax(<8 x float> %in)596 ret float %res597}598 599define float @reduce_fmax_float_reassoc_nonpow2(<7 x float> %in) {600; CHECK-SM80-LABEL: reduce_fmax_float_reassoc_nonpow2(601; CHECK-SM80: {602; CHECK-SM80-NEXT: .reg .b32 %r<14>;603; CHECK-SM80-EMPTY:604; CHECK-SM80-NEXT: // %bb.0:605; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmax_float_reassoc_nonpow2_param_0+24];606; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmax_float_reassoc_nonpow2_param_0+16];607; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_reassoc_nonpow2_param_0];608; CHECK-SM80-NEXT: max.f32 %r8, %r5, %r6;609; CHECK-SM80-NEXT: max.f32 %r9, %r8, %r7;610; CHECK-SM80-NEXT: max.f32 %r10, %r3, %r4;611; CHECK-SM80-NEXT: max.f32 %r11, %r1, %r2;612; CHECK-SM80-NEXT: max.f32 %r12, %r11, %r10;613; CHECK-SM80-NEXT: max.f32 %r13, %r12, %r9;614; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;615; CHECK-SM80-NEXT: ret;616;617; CHECK-SM100-LABEL: reduce_fmax_float_reassoc_nonpow2(618; CHECK-SM100: {619; CHECK-SM100-NEXT: .reg .b32 %r<11>;620; CHECK-SM100-EMPTY:621; CHECK-SM100-NEXT: // %bb.0:622; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmax_float_reassoc_nonpow2_param_0+24];623; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmax_float_reassoc_nonpow2_param_0+16];624; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_reassoc_nonpow2_param_0];625; CHECK-SM100-NEXT: max.f32 %r8, %r4, %r5, %r6;626; CHECK-SM100-NEXT: max.f32 %r9, %r1, %r2, %r3;627; CHECK-SM100-NEXT: max.f32 %r10, %r9, %r8, %r7;628; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;629; CHECK-SM100-NEXT: ret;630 %res = call reassoc float @llvm.vector.reduce.fmax(<7 x float> %in)631 ret float %res632}633 634define float @reduce_fmax_float_nnan(<8 x float> %in) {635; CHECK-SM80-LABEL: reduce_fmax_float_nnan(636; CHECK-SM80: {637; CHECK-SM80-NEXT: .reg .b32 %r<16>;638; CHECK-SM80-EMPTY:639; CHECK-SM80-NEXT: // %bb.0:640; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_nnan_param_0+16];641; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_nnan_param_0];642; CHECK-SM80-NEXT: max.f32 %r9, %r7, %r8;643; CHECK-SM80-NEXT: max.f32 %r10, %r5, %r6;644; CHECK-SM80-NEXT: max.f32 %r11, %r10, %r9;645; CHECK-SM80-NEXT: max.f32 %r12, %r3, %r4;646; CHECK-SM80-NEXT: max.f32 %r13, %r1, %r2;647; CHECK-SM80-NEXT: max.f32 %r14, %r13, %r12;648; CHECK-SM80-NEXT: max.f32 %r15, %r14, %r11;649; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;650; CHECK-SM80-NEXT: ret;651;652; CHECK-SM100-LABEL: reduce_fmax_float_nnan(653; CHECK-SM100: {654; CHECK-SM100-NEXT: .reg .b32 %r<13>;655; CHECK-SM100-NEXT: .reg .b64 %rd<5>;656; CHECK-SM100-EMPTY:657; CHECK-SM100-NEXT: // %bb.0:658; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmax_float_nnan_param_0+16];659; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;660; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmax_float_nnan_param_0];661; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;662; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;663; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;664; CHECK-SM100-NEXT: max.f32 %r9, %r8, %r5, %r6;665; CHECK-SM100-NEXT: max.f32 %r10, %r3, %r4, %r7;666; CHECK-SM100-NEXT: max.f32 %r11, %r10, %r9, %r1;667; CHECK-SM100-NEXT: max.f32 %r12, %r11, %r2;668; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;669; CHECK-SM100-NEXT: ret;670 %res = call nnan float @llvm.vector.reduce.fmax(<8 x float> %in)671 ret float %res672}673 674define float @reduce_fmax_float_nnan_nonpow2(<7 x float> %in) {675; CHECK-SM80-LABEL: reduce_fmax_float_nnan_nonpow2(676; CHECK-SM80: {677; CHECK-SM80-NEXT: .reg .b32 %r<14>;678; CHECK-SM80-EMPTY:679; CHECK-SM80-NEXT: // %bb.0:680; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmax_float_nnan_nonpow2_param_0+24];681; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmax_float_nnan_nonpow2_param_0+16];682; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_nnan_nonpow2_param_0];683; CHECK-SM80-NEXT: max.f32 %r8, %r5, %r6;684; CHECK-SM80-NEXT: max.f32 %r9, %r8, %r7;685; CHECK-SM80-NEXT: max.f32 %r10, %r3, %r4;686; CHECK-SM80-NEXT: max.f32 %r11, %r1, %r2;687; CHECK-SM80-NEXT: max.f32 %r12, %r11, %r10;688; CHECK-SM80-NEXT: max.f32 %r13, %r12, %r9;689; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;690; CHECK-SM80-NEXT: ret;691;692; CHECK-SM100-LABEL: reduce_fmax_float_nnan_nonpow2(693; CHECK-SM100: {694; CHECK-SM100-NEXT: .reg .b32 %r<11>;695; CHECK-SM100-EMPTY:696; CHECK-SM100-NEXT: // %bb.0:697; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmax_float_nnan_nonpow2_param_0+24];698; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmax_float_nnan_nonpow2_param_0+16];699; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_nnan_nonpow2_param_0];700; CHECK-SM100-NEXT: max.f32 %r8, %r4, %r5, %r6;701; CHECK-SM100-NEXT: max.f32 %r9, %r1, %r2, %r3;702; CHECK-SM100-NEXT: max.f32 %r10, %r9, %r8, %r7;703; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;704; CHECK-SM100-NEXT: ret;705 %res = call nnan float @llvm.vector.reduce.fmax(<7 x float> %in)706 ret float %res707}708 709define half @reduce_fmin_half(<8 x half> %in) {710; CHECK-LABEL: reduce_fmin_half(711; CHECK: {712; CHECK-NEXT: .reg .b16 %rs<4>;713; CHECK-NEXT: .reg .b32 %r<8>;714; CHECK-EMPTY:715; CHECK-NEXT: // %bb.0:716; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_half_param_0];717; CHECK-NEXT: min.f16x2 %r5, %r2, %r4;718; CHECK-NEXT: min.f16x2 %r6, %r1, %r3;719; CHECK-NEXT: min.f16x2 %r7, %r6, %r5;720; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;721; CHECK-NEXT: min.f16 %rs3, %rs1, %rs2;722; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;723; CHECK-NEXT: ret;724 %res = call half @llvm.vector.reduce.fmin(<8 x half> %in)725 ret half %res726}727 728define half @reduce_fmin_half_reassoc(<8 x half> %in) {729; CHECK-LABEL: reduce_fmin_half_reassoc(730; CHECK: {731; CHECK-NEXT: .reg .b16 %rs<4>;732; CHECK-NEXT: .reg .b32 %r<8>;733; CHECK-EMPTY:734; CHECK-NEXT: // %bb.0:735; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_half_reassoc_param_0];736; CHECK-NEXT: min.f16x2 %r5, %r2, %r4;737; CHECK-NEXT: min.f16x2 %r6, %r1, %r3;738; CHECK-NEXT: min.f16x2 %r7, %r6, %r5;739; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;740; CHECK-NEXT: min.f16 %rs3, %rs1, %rs2;741; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;742; CHECK-NEXT: ret;743 %res = call reassoc half @llvm.vector.reduce.fmin(<8 x half> %in)744 ret half %res745}746 747define half @reduce_fmin_half_reassoc_nonpow2(<7 x half> %in) {748; CHECK-LABEL: reduce_fmin_half_reassoc_nonpow2(749; CHECK: {750; CHECK-NEXT: .reg .b16 %rs<12>;751; CHECK-NEXT: .reg .b32 %r<8>;752; CHECK-EMPTY:753; CHECK-NEXT: // %bb.0:754; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmin_half_reassoc_nonpow2_param_0+8];755; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;756; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmin_half_reassoc_nonpow2_param_0];757; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;758; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;759; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmin_half_reassoc_nonpow2_param_0+12];760; CHECK-NEXT: min.f16x2 %r4, %r2, %r1;761; CHECK-NEXT: mov.b16 %rs8, 0x7E00;762; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};763; CHECK-NEXT: min.f16x2 %r6, %r3, %r5;764; CHECK-NEXT: min.f16x2 %r7, %r4, %r6;765; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;766; CHECK-NEXT: min.f16 %rs11, %rs9, %rs10;767; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;768; CHECK-NEXT: ret;769 %res = call reassoc half @llvm.vector.reduce.fmin(<7 x half> %in)770 ret half %res771}772 773define half @reduce_fmin_half_nnan(<8 x half> %in) {774; CHECK-LABEL: reduce_fmin_half_nnan(775; CHECK: {776; CHECK-NEXT: .reg .b16 %rs<4>;777; CHECK-NEXT: .reg .b32 %r<8>;778; CHECK-EMPTY:779; CHECK-NEXT: // %bb.0:780; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_half_nnan_param_0];781; CHECK-NEXT: min.f16x2 %r5, %r2, %r4;782; CHECK-NEXT: min.f16x2 %r6, %r1, %r3;783; CHECK-NEXT: min.f16x2 %r7, %r6, %r5;784; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;785; CHECK-NEXT: min.f16 %rs3, %rs1, %rs2;786; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;787; CHECK-NEXT: ret;788 %res = call nnan half @llvm.vector.reduce.fmin(<8 x half> %in)789 ret half %res790}791 792define half @reduce_fmin_half_nnan_nonpow2(<7 x half> %in) {793; CHECK-LABEL: reduce_fmin_half_nnan_nonpow2(794; CHECK: {795; CHECK-NEXT: .reg .b16 %rs<12>;796; CHECK-NEXT: .reg .b32 %r<8>;797; CHECK-EMPTY:798; CHECK-NEXT: // %bb.0:799; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmin_half_nnan_nonpow2_param_0+8];800; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;801; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmin_half_nnan_nonpow2_param_0];802; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;803; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;804; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmin_half_nnan_nonpow2_param_0+12];805; CHECK-NEXT: min.f16x2 %r4, %r2, %r1;806; CHECK-NEXT: mov.b16 %rs8, 0x7C00;807; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};808; CHECK-NEXT: min.f16x2 %r6, %r3, %r5;809; CHECK-NEXT: min.f16x2 %r7, %r4, %r6;810; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;811; CHECK-NEXT: min.f16 %rs11, %rs9, %rs10;812; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;813; CHECK-NEXT: ret;814 %res = call nnan half @llvm.vector.reduce.fmin(<7 x half> %in)815 ret half %res816}817 818define float @reduce_fmin_float(<8 x float> %in) {819; CHECK-SM80-LABEL: reduce_fmin_float(820; CHECK-SM80: {821; CHECK-SM80-NEXT: .reg .b32 %r<16>;822; CHECK-SM80-EMPTY:823; CHECK-SM80-NEXT: // %bb.0:824; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_param_0+16];825; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_param_0];826; CHECK-SM80-NEXT: min.f32 %r9, %r7, %r8;827; CHECK-SM80-NEXT: min.f32 %r10, %r5, %r6;828; CHECK-SM80-NEXT: min.f32 %r11, %r10, %r9;829; CHECK-SM80-NEXT: min.f32 %r12, %r3, %r4;830; CHECK-SM80-NEXT: min.f32 %r13, %r1, %r2;831; CHECK-SM80-NEXT: min.f32 %r14, %r13, %r12;832; CHECK-SM80-NEXT: min.f32 %r15, %r14, %r11;833; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;834; CHECK-SM80-NEXT: ret;835;836; CHECK-SM100-LABEL: reduce_fmin_float(837; CHECK-SM100: {838; CHECK-SM100-NEXT: .reg .b32 %r<13>;839; CHECK-SM100-NEXT: .reg .b64 %rd<5>;840; CHECK-SM100-EMPTY:841; CHECK-SM100-NEXT: // %bb.0:842; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmin_float_param_0+16];843; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;844; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmin_float_param_0];845; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;846; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;847; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;848; CHECK-SM100-NEXT: min.f32 %r9, %r8, %r5, %r6;849; CHECK-SM100-NEXT: min.f32 %r10, %r3, %r4, %r7;850; CHECK-SM100-NEXT: min.f32 %r11, %r10, %r9, %r1;851; CHECK-SM100-NEXT: min.f32 %r12, %r11, %r2;852; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;853; CHECK-SM100-NEXT: ret;854 %res = call float @llvm.vector.reduce.fmin(<8 x float> %in)855 ret float %res856}857 858define float @reduce_fmin_float_reassoc(<8 x float> %in) {859; CHECK-SM80-LABEL: reduce_fmin_float_reassoc(860; CHECK-SM80: {861; CHECK-SM80-NEXT: .reg .b32 %r<16>;862; CHECK-SM80-EMPTY:863; CHECK-SM80-NEXT: // %bb.0:864; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_reassoc_param_0+16];865; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_reassoc_param_0];866; CHECK-SM80-NEXT: min.f32 %r9, %r7, %r8;867; CHECK-SM80-NEXT: min.f32 %r10, %r5, %r6;868; CHECK-SM80-NEXT: min.f32 %r11, %r10, %r9;869; CHECK-SM80-NEXT: min.f32 %r12, %r3, %r4;870; CHECK-SM80-NEXT: min.f32 %r13, %r1, %r2;871; CHECK-SM80-NEXT: min.f32 %r14, %r13, %r12;872; CHECK-SM80-NEXT: min.f32 %r15, %r14, %r11;873; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;874; CHECK-SM80-NEXT: ret;875;876; CHECK-SM100-LABEL: reduce_fmin_float_reassoc(877; CHECK-SM100: {878; CHECK-SM100-NEXT: .reg .b32 %r<13>;879; CHECK-SM100-NEXT: .reg .b64 %rd<5>;880; CHECK-SM100-EMPTY:881; CHECK-SM100-NEXT: // %bb.0:882; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmin_float_reassoc_param_0+16];883; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;884; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmin_float_reassoc_param_0];885; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;886; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;887; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;888; CHECK-SM100-NEXT: min.f32 %r9, %r8, %r5, %r6;889; CHECK-SM100-NEXT: min.f32 %r10, %r3, %r4, %r7;890; CHECK-SM100-NEXT: min.f32 %r11, %r10, %r9, %r1;891; CHECK-SM100-NEXT: min.f32 %r12, %r11, %r2;892; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;893; CHECK-SM100-NEXT: ret;894 %res = call reassoc float @llvm.vector.reduce.fmin(<8 x float> %in)895 ret float %res896}897 898define float @reduce_fmin_float_reassoc_nonpow2(<7 x float> %in) {899; CHECK-SM80-LABEL: reduce_fmin_float_reassoc_nonpow2(900; CHECK-SM80: {901; CHECK-SM80-NEXT: .reg .b32 %r<14>;902; CHECK-SM80-EMPTY:903; CHECK-SM80-NEXT: // %bb.0:904; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmin_float_reassoc_nonpow2_param_0+24];905; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmin_float_reassoc_nonpow2_param_0+16];906; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_reassoc_nonpow2_param_0];907; CHECK-SM80-NEXT: min.f32 %r8, %r5, %r6;908; CHECK-SM80-NEXT: min.f32 %r9, %r8, %r7;909; CHECK-SM80-NEXT: min.f32 %r10, %r3, %r4;910; CHECK-SM80-NEXT: min.f32 %r11, %r1, %r2;911; CHECK-SM80-NEXT: min.f32 %r12, %r11, %r10;912; CHECK-SM80-NEXT: min.f32 %r13, %r12, %r9;913; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;914; CHECK-SM80-NEXT: ret;915;916; CHECK-SM100-LABEL: reduce_fmin_float_reassoc_nonpow2(917; CHECK-SM100: {918; CHECK-SM100-NEXT: .reg .b32 %r<11>;919; CHECK-SM100-EMPTY:920; CHECK-SM100-NEXT: // %bb.0:921; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmin_float_reassoc_nonpow2_param_0+24];922; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmin_float_reassoc_nonpow2_param_0+16];923; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_reassoc_nonpow2_param_0];924; CHECK-SM100-NEXT: min.f32 %r8, %r4, %r5, %r6;925; CHECK-SM100-NEXT: min.f32 %r9, %r1, %r2, %r3;926; CHECK-SM100-NEXT: min.f32 %r10, %r9, %r8, %r7;927; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;928; CHECK-SM100-NEXT: ret;929 %res = call reassoc float @llvm.vector.reduce.fmin(<7 x float> %in)930 ret float %res931}932 933define float @reduce_fmin_float_nnan(<8 x float> %in) {934; CHECK-SM80-LABEL: reduce_fmin_float_nnan(935; CHECK-SM80: {936; CHECK-SM80-NEXT: .reg .b32 %r<16>;937; CHECK-SM80-EMPTY:938; CHECK-SM80-NEXT: // %bb.0:939; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_nnan_param_0+16];940; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_nnan_param_0];941; CHECK-SM80-NEXT: min.f32 %r9, %r7, %r8;942; CHECK-SM80-NEXT: min.f32 %r10, %r5, %r6;943; CHECK-SM80-NEXT: min.f32 %r11, %r10, %r9;944; CHECK-SM80-NEXT: min.f32 %r12, %r3, %r4;945; CHECK-SM80-NEXT: min.f32 %r13, %r1, %r2;946; CHECK-SM80-NEXT: min.f32 %r14, %r13, %r12;947; CHECK-SM80-NEXT: min.f32 %r15, %r14, %r11;948; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;949; CHECK-SM80-NEXT: ret;950;951; CHECK-SM100-LABEL: reduce_fmin_float_nnan(952; CHECK-SM100: {953; CHECK-SM100-NEXT: .reg .b32 %r<13>;954; CHECK-SM100-NEXT: .reg .b64 %rd<5>;955; CHECK-SM100-EMPTY:956; CHECK-SM100-NEXT: // %bb.0:957; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmin_float_nnan_param_0+16];958; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;959; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmin_float_nnan_param_0];960; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;961; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;962; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;963; CHECK-SM100-NEXT: min.f32 %r9, %r8, %r5, %r6;964; CHECK-SM100-NEXT: min.f32 %r10, %r3, %r4, %r7;965; CHECK-SM100-NEXT: min.f32 %r11, %r10, %r9, %r1;966; CHECK-SM100-NEXT: min.f32 %r12, %r11, %r2;967; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;968; CHECK-SM100-NEXT: ret;969 %res = call nnan float @llvm.vector.reduce.fmin(<8 x float> %in)970 ret float %res971}972 973define float @reduce_fmin_float_nnan_nonpow2(<7 x float> %in) {974; CHECK-SM80-LABEL: reduce_fmin_float_nnan_nonpow2(975; CHECK-SM80: {976; CHECK-SM80-NEXT: .reg .b32 %r<14>;977; CHECK-SM80-EMPTY:978; CHECK-SM80-NEXT: // %bb.0:979; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmin_float_nnan_nonpow2_param_0+24];980; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmin_float_nnan_nonpow2_param_0+16];981; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_nnan_nonpow2_param_0];982; CHECK-SM80-NEXT: min.f32 %r8, %r5, %r6;983; CHECK-SM80-NEXT: min.f32 %r9, %r8, %r7;984; CHECK-SM80-NEXT: min.f32 %r10, %r3, %r4;985; CHECK-SM80-NEXT: min.f32 %r11, %r1, %r2;986; CHECK-SM80-NEXT: min.f32 %r12, %r11, %r10;987; CHECK-SM80-NEXT: min.f32 %r13, %r12, %r9;988; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;989; CHECK-SM80-NEXT: ret;990;991; CHECK-SM100-LABEL: reduce_fmin_float_nnan_nonpow2(992; CHECK-SM100: {993; CHECK-SM100-NEXT: .reg .b32 %r<11>;994; CHECK-SM100-EMPTY:995; CHECK-SM100-NEXT: // %bb.0:996; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmin_float_nnan_nonpow2_param_0+24];997; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmin_float_nnan_nonpow2_param_0+16];998; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_nnan_nonpow2_param_0];999; CHECK-SM100-NEXT: min.f32 %r8, %r4, %r5, %r6;1000; CHECK-SM100-NEXT: min.f32 %r9, %r1, %r2, %r3;1001; CHECK-SM100-NEXT: min.f32 %r10, %r9, %r8, %r7;1002; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;1003; CHECK-SM100-NEXT: ret;1004 %res = call nnan float @llvm.vector.reduce.fmin(<7 x float> %in)1005 ret float %res1006}1007 1008define half @reduce_fmaximum_half(<8 x half> %in) {1009; CHECK-LABEL: reduce_fmaximum_half(1010; CHECK: {1011; CHECK-NEXT: .reg .b16 %rs<4>;1012; CHECK-NEXT: .reg .b32 %r<8>;1013; CHECK-EMPTY:1014; CHECK-NEXT: // %bb.0:1015; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_half_param_0];1016; CHECK-NEXT: max.NaN.f16x2 %r5, %r2, %r4;1017; CHECK-NEXT: max.NaN.f16x2 %r6, %r1, %r3;1018; CHECK-NEXT: max.NaN.f16x2 %r7, %r6, %r5;1019; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;1020; CHECK-NEXT: max.NaN.f16 %rs3, %rs1, %rs2;1021; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;1022; CHECK-NEXT: ret;1023 %res = call half @llvm.vector.reduce.fmaximum(<8 x half> %in)1024 ret half %res1025}1026 1027define half @reduce_fmaximum_half_reassoc(<8 x half> %in) {1028; CHECK-LABEL: reduce_fmaximum_half_reassoc(1029; CHECK: {1030; CHECK-NEXT: .reg .b16 %rs<4>;1031; CHECK-NEXT: .reg .b32 %r<8>;1032; CHECK-EMPTY:1033; CHECK-NEXT: // %bb.0:1034; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_half_reassoc_param_0];1035; CHECK-NEXT: max.NaN.f16x2 %r5, %r2, %r4;1036; CHECK-NEXT: max.NaN.f16x2 %r6, %r1, %r3;1037; CHECK-NEXT: max.NaN.f16x2 %r7, %r6, %r5;1038; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;1039; CHECK-NEXT: max.NaN.f16 %rs3, %rs1, %rs2;1040; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;1041; CHECK-NEXT: ret;1042 %res = call reassoc half @llvm.vector.reduce.fmaximum(<8 x half> %in)1043 ret half %res1044}1045 1046define half @reduce_fmaximum_half_reassoc_nonpow2(<7 x half> %in) {1047; CHECK-LABEL: reduce_fmaximum_half_reassoc_nonpow2(1048; CHECK: {1049; CHECK-NEXT: .reg .b16 %rs<12>;1050; CHECK-NEXT: .reg .b32 %r<8>;1051; CHECK-EMPTY:1052; CHECK-NEXT: // %bb.0:1053; CHECK-NEXT: ld.param.b32 %r1, [reduce_fmaximum_half_reassoc_nonpow2_param_0+8];1054; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;1055; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fmaximum_half_reassoc_nonpow2_param_0];1056; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;1057; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;1058; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fmaximum_half_reassoc_nonpow2_param_0+12];1059; CHECK-NEXT: max.NaN.f16x2 %r4, %r2, %r1;1060; CHECK-NEXT: mov.b16 %rs8, 0xFC00;1061; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};1062; CHECK-NEXT: max.NaN.f16x2 %r6, %r3, %r5;1063; CHECK-NEXT: max.NaN.f16x2 %r7, %r4, %r6;1064; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;1065; CHECK-NEXT: max.NaN.f16 %rs11, %rs9, %rs10;1066; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;1067; CHECK-NEXT: ret;1068 %res = call reassoc half @llvm.vector.reduce.fmaximum(<7 x half> %in)1069 ret half %res1070}1071 1072define float @reduce_fmaximum_float(<8 x float> %in) {1073; CHECK-SM80-LABEL: reduce_fmaximum_float(1074; CHECK-SM80: {1075; CHECK-SM80-NEXT: .reg .b32 %r<16>;1076; CHECK-SM80-EMPTY:1077; CHECK-SM80-NEXT: // %bb.0:1078; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_param_0+16];1079; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_param_0];1080; CHECK-SM80-NEXT: max.NaN.f32 %r9, %r7, %r8;1081; CHECK-SM80-NEXT: max.NaN.f32 %r10, %r5, %r6;1082; CHECK-SM80-NEXT: max.NaN.f32 %r11, %r10, %r9;1083; CHECK-SM80-NEXT: max.NaN.f32 %r12, %r3, %r4;1084; CHECK-SM80-NEXT: max.NaN.f32 %r13, %r1, %r2;1085; CHECK-SM80-NEXT: max.NaN.f32 %r14, %r13, %r12;1086; CHECK-SM80-NEXT: max.NaN.f32 %r15, %r14, %r11;1087; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1088; CHECK-SM80-NEXT: ret;1089;1090; CHECK-SM100-LABEL: reduce_fmaximum_float(1091; CHECK-SM100: {1092; CHECK-SM100-NEXT: .reg .b32 %r<13>;1093; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1094; CHECK-SM100-EMPTY:1095; CHECK-SM100-NEXT: // %bb.0:1096; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmaximum_float_param_0+16];1097; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1098; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmaximum_float_param_0];1099; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;1100; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;1101; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;1102; CHECK-SM100-NEXT: max.NaN.f32 %r9, %r8, %r5, %r6;1103; CHECK-SM100-NEXT: max.NaN.f32 %r10, %r3, %r4, %r7;1104; CHECK-SM100-NEXT: max.NaN.f32 %r11, %r10, %r9, %r1;1105; CHECK-SM100-NEXT: max.NaN.f32 %r12, %r11, %r2;1106; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;1107; CHECK-SM100-NEXT: ret;1108 %res = call float @llvm.vector.reduce.fmaximum(<8 x float> %in)1109 ret float %res1110}1111 1112define float @reduce_fmaximum_float_reassoc(<8 x float> %in) {1113; CHECK-SM80-LABEL: reduce_fmaximum_float_reassoc(1114; CHECK-SM80: {1115; CHECK-SM80-NEXT: .reg .b32 %r<16>;1116; CHECK-SM80-EMPTY:1117; CHECK-SM80-NEXT: // %bb.0:1118; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_reassoc_param_0+16];1119; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_reassoc_param_0];1120; CHECK-SM80-NEXT: max.NaN.f32 %r9, %r7, %r8;1121; CHECK-SM80-NEXT: max.NaN.f32 %r10, %r5, %r6;1122; CHECK-SM80-NEXT: max.NaN.f32 %r11, %r10, %r9;1123; CHECK-SM80-NEXT: max.NaN.f32 %r12, %r3, %r4;1124; CHECK-SM80-NEXT: max.NaN.f32 %r13, %r1, %r2;1125; CHECK-SM80-NEXT: max.NaN.f32 %r14, %r13, %r12;1126; CHECK-SM80-NEXT: max.NaN.f32 %r15, %r14, %r11;1127; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1128; CHECK-SM80-NEXT: ret;1129;1130; CHECK-SM100-LABEL: reduce_fmaximum_float_reassoc(1131; CHECK-SM100: {1132; CHECK-SM100-NEXT: .reg .b32 %r<13>;1133; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1134; CHECK-SM100-EMPTY:1135; CHECK-SM100-NEXT: // %bb.0:1136; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmaximum_float_reassoc_param_0+16];1137; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1138; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmaximum_float_reassoc_param_0];1139; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;1140; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;1141; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;1142; CHECK-SM100-NEXT: max.NaN.f32 %r9, %r8, %r5, %r6;1143; CHECK-SM100-NEXT: max.NaN.f32 %r10, %r3, %r4, %r7;1144; CHECK-SM100-NEXT: max.NaN.f32 %r11, %r10, %r9, %r1;1145; CHECK-SM100-NEXT: max.NaN.f32 %r12, %r11, %r2;1146; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;1147; CHECK-SM100-NEXT: ret;1148 %res = call reassoc float @llvm.vector.reduce.fmaximum(<8 x float> %in)1149 ret float %res1150}1151 1152define float @reduce_fmaximum_float_reassoc_nonpow2(<7 x float> %in) {1153; CHECK-SM80-LABEL: reduce_fmaximum_float_reassoc_nonpow2(1154; CHECK-SM80: {1155; CHECK-SM80-NEXT: .reg .b32 %r<14>;1156; CHECK-SM80-EMPTY:1157; CHECK-SM80-NEXT: // %bb.0:1158; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fmaximum_float_reassoc_nonpow2_param_0+24];1159; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmaximum_float_reassoc_nonpow2_param_0+16];1160; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_reassoc_nonpow2_param_0];1161; CHECK-SM80-NEXT: max.NaN.f32 %r8, %r5, %r6;1162; CHECK-SM80-NEXT: max.NaN.f32 %r9, %r8, %r7;1163; CHECK-SM80-NEXT: max.NaN.f32 %r10, %r3, %r4;1164; CHECK-SM80-NEXT: max.NaN.f32 %r11, %r1, %r2;1165; CHECK-SM80-NEXT: max.NaN.f32 %r12, %r11, %r10;1166; CHECK-SM80-NEXT: max.NaN.f32 %r13, %r12, %r9;1167; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;1168; CHECK-SM80-NEXT: ret;1169;1170; CHECK-SM100-LABEL: reduce_fmaximum_float_reassoc_nonpow2(1171; CHECK-SM100: {1172; CHECK-SM100-NEXT: .reg .b32 %r<11>;1173; CHECK-SM100-EMPTY:1174; CHECK-SM100-NEXT: // %bb.0:1175; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fmaximum_float_reassoc_nonpow2_param_0+24];1176; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fmaximum_float_reassoc_nonpow2_param_0+16];1177; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_reassoc_nonpow2_param_0];1178; CHECK-SM100-NEXT: max.NaN.f32 %r8, %r4, %r5, %r6;1179; CHECK-SM100-NEXT: max.NaN.f32 %r9, %r1, %r2, %r3;1180; CHECK-SM100-NEXT: max.NaN.f32 %r10, %r9, %r8, %r7;1181; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;1182; CHECK-SM100-NEXT: ret;1183 %res = call reassoc float @llvm.vector.reduce.fmaximum(<7 x float> %in)1184 ret float %res1185}1186 1187define half @reduce_fminimum_half(<8 x half> %in) {1188; CHECK-LABEL: reduce_fminimum_half(1189; CHECK: {1190; CHECK-NEXT: .reg .b16 %rs<4>;1191; CHECK-NEXT: .reg .b32 %r<8>;1192; CHECK-EMPTY:1193; CHECK-NEXT: // %bb.0:1194; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_half_param_0];1195; CHECK-NEXT: min.NaN.f16x2 %r5, %r2, %r4;1196; CHECK-NEXT: min.NaN.f16x2 %r6, %r1, %r3;1197; CHECK-NEXT: min.NaN.f16x2 %r7, %r6, %r5;1198; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;1199; CHECK-NEXT: min.NaN.f16 %rs3, %rs1, %rs2;1200; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;1201; CHECK-NEXT: ret;1202 %res = call half @llvm.vector.reduce.fminimum(<8 x half> %in)1203 ret half %res1204}1205 1206define half @reduce_fminimum_half_reassoc(<8 x half> %in) {1207; CHECK-LABEL: reduce_fminimum_half_reassoc(1208; CHECK: {1209; CHECK-NEXT: .reg .b16 %rs<4>;1210; CHECK-NEXT: .reg .b32 %r<8>;1211; CHECK-EMPTY:1212; CHECK-NEXT: // %bb.0:1213; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_half_reassoc_param_0];1214; CHECK-NEXT: min.NaN.f16x2 %r5, %r2, %r4;1215; CHECK-NEXT: min.NaN.f16x2 %r6, %r1, %r3;1216; CHECK-NEXT: min.NaN.f16x2 %r7, %r6, %r5;1217; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;1218; CHECK-NEXT: min.NaN.f16 %rs3, %rs1, %rs2;1219; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;1220; CHECK-NEXT: ret;1221 %res = call reassoc half @llvm.vector.reduce.fminimum(<8 x half> %in)1222 ret half %res1223}1224 1225define half @reduce_fminimum_half_reassoc_nonpow2(<7 x half> %in) {1226; CHECK-LABEL: reduce_fminimum_half_reassoc_nonpow2(1227; CHECK: {1228; CHECK-NEXT: .reg .b16 %rs<12>;1229; CHECK-NEXT: .reg .b32 %r<8>;1230; CHECK-EMPTY:1231; CHECK-NEXT: // %bb.0:1232; CHECK-NEXT: ld.param.b32 %r1, [reduce_fminimum_half_reassoc_nonpow2_param_0+8];1233; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;1234; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_fminimum_half_reassoc_nonpow2_param_0];1235; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;1236; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;1237; CHECK-NEXT: ld.param.b16 %rs7, [reduce_fminimum_half_reassoc_nonpow2_param_0+12];1238; CHECK-NEXT: min.NaN.f16x2 %r4, %r2, %r1;1239; CHECK-NEXT: mov.b16 %rs8, 0x7C00;1240; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};1241; CHECK-NEXT: min.NaN.f16x2 %r6, %r3, %r5;1242; CHECK-NEXT: min.NaN.f16x2 %r7, %r4, %r6;1243; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;1244; CHECK-NEXT: min.NaN.f16 %rs11, %rs9, %rs10;1245; CHECK-NEXT: st.param.b16 [func_retval0], %rs11;1246; CHECK-NEXT: ret;1247 %res = call reassoc half @llvm.vector.reduce.fminimum(<7 x half> %in)1248 ret half %res1249}1250 1251define float @reduce_fminimum_float(<8 x float> %in) {1252; CHECK-SM80-LABEL: reduce_fminimum_float(1253; CHECK-SM80: {1254; CHECK-SM80-NEXT: .reg .b32 %r<16>;1255; CHECK-SM80-EMPTY:1256; CHECK-SM80-NEXT: // %bb.0:1257; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_param_0+16];1258; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_param_0];1259; CHECK-SM80-NEXT: min.NaN.f32 %r9, %r7, %r8;1260; CHECK-SM80-NEXT: min.NaN.f32 %r10, %r5, %r6;1261; CHECK-SM80-NEXT: min.NaN.f32 %r11, %r10, %r9;1262; CHECK-SM80-NEXT: min.NaN.f32 %r12, %r3, %r4;1263; CHECK-SM80-NEXT: min.NaN.f32 %r13, %r1, %r2;1264; CHECK-SM80-NEXT: min.NaN.f32 %r14, %r13, %r12;1265; CHECK-SM80-NEXT: min.NaN.f32 %r15, %r14, %r11;1266; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1267; CHECK-SM80-NEXT: ret;1268;1269; CHECK-SM100-LABEL: reduce_fminimum_float(1270; CHECK-SM100: {1271; CHECK-SM100-NEXT: .reg .b32 %r<13>;1272; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1273; CHECK-SM100-EMPTY:1274; CHECK-SM100-NEXT: // %bb.0:1275; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fminimum_float_param_0+16];1276; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1277; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fminimum_float_param_0];1278; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;1279; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;1280; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;1281; CHECK-SM100-NEXT: min.NaN.f32 %r9, %r8, %r5, %r6;1282; CHECK-SM100-NEXT: min.NaN.f32 %r10, %r3, %r4, %r7;1283; CHECK-SM100-NEXT: min.NaN.f32 %r11, %r10, %r9, %r1;1284; CHECK-SM100-NEXT: min.NaN.f32 %r12, %r11, %r2;1285; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;1286; CHECK-SM100-NEXT: ret;1287 %res = call float @llvm.vector.reduce.fminimum(<8 x float> %in)1288 ret float %res1289}1290 1291define float @reduce_fminimum_float_reassoc(<8 x float> %in) {1292; CHECK-SM80-LABEL: reduce_fminimum_float_reassoc(1293; CHECK-SM80: {1294; CHECK-SM80-NEXT: .reg .b32 %r<16>;1295; CHECK-SM80-EMPTY:1296; CHECK-SM80-NEXT: // %bb.0:1297; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_reassoc_param_0+16];1298; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_reassoc_param_0];1299; CHECK-SM80-NEXT: min.NaN.f32 %r9, %r7, %r8;1300; CHECK-SM80-NEXT: min.NaN.f32 %r10, %r5, %r6;1301; CHECK-SM80-NEXT: min.NaN.f32 %r11, %r10, %r9;1302; CHECK-SM80-NEXT: min.NaN.f32 %r12, %r3, %r4;1303; CHECK-SM80-NEXT: min.NaN.f32 %r13, %r1, %r2;1304; CHECK-SM80-NEXT: min.NaN.f32 %r14, %r13, %r12;1305; CHECK-SM80-NEXT: min.NaN.f32 %r15, %r14, %r11;1306; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1307; CHECK-SM80-NEXT: ret;1308;1309; CHECK-SM100-LABEL: reduce_fminimum_float_reassoc(1310; CHECK-SM100: {1311; CHECK-SM100-NEXT: .reg .b32 %r<13>;1312; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1313; CHECK-SM100-EMPTY:1314; CHECK-SM100-NEXT: // %bb.0:1315; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fminimum_float_reassoc_param_0+16];1316; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1317; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fminimum_float_reassoc_param_0];1318; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd1;1319; CHECK-SM100-NEXT: mov.b64 {%r5, %r6}, %rd3;1320; CHECK-SM100-NEXT: mov.b64 {%r7, %r8}, %rd2;1321; CHECK-SM100-NEXT: min.NaN.f32 %r9, %r8, %r5, %r6;1322; CHECK-SM100-NEXT: min.NaN.f32 %r10, %r3, %r4, %r7;1323; CHECK-SM100-NEXT: min.NaN.f32 %r11, %r10, %r9, %r1;1324; CHECK-SM100-NEXT: min.NaN.f32 %r12, %r11, %r2;1325; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r12;1326; CHECK-SM100-NEXT: ret;1327 %res = call reassoc float @llvm.vector.reduce.fminimum(<8 x float> %in)1328 ret float %res1329}1330 1331define float @reduce_fminimum_float_reassoc_nonpow2(<7 x float> %in) {1332; CHECK-SM80-LABEL: reduce_fminimum_float_reassoc_nonpow2(1333; CHECK-SM80: {1334; CHECK-SM80-NEXT: .reg .b32 %r<14>;1335; CHECK-SM80-EMPTY:1336; CHECK-SM80-NEXT: // %bb.0:1337; CHECK-SM80-NEXT: ld.param.b32 %r7, [reduce_fminimum_float_reassoc_nonpow2_param_0+24];1338; CHECK-SM80-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fminimum_float_reassoc_nonpow2_param_0+16];1339; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_reassoc_nonpow2_param_0];1340; CHECK-SM80-NEXT: min.NaN.f32 %r8, %r5, %r6;1341; CHECK-SM80-NEXT: min.NaN.f32 %r9, %r8, %r7;1342; CHECK-SM80-NEXT: min.NaN.f32 %r10, %r3, %r4;1343; CHECK-SM80-NEXT: min.NaN.f32 %r11, %r1, %r2;1344; CHECK-SM80-NEXT: min.NaN.f32 %r12, %r11, %r10;1345; CHECK-SM80-NEXT: min.NaN.f32 %r13, %r12, %r9;1346; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r13;1347; CHECK-SM80-NEXT: ret;1348;1349; CHECK-SM100-LABEL: reduce_fminimum_float_reassoc_nonpow2(1350; CHECK-SM100: {1351; CHECK-SM100-NEXT: .reg .b32 %r<11>;1352; CHECK-SM100-EMPTY:1353; CHECK-SM100-NEXT: // %bb.0:1354; CHECK-SM100-NEXT: ld.param.b32 %r7, [reduce_fminimum_float_reassoc_nonpow2_param_0+24];1355; CHECK-SM100-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_fminimum_float_reassoc_nonpow2_param_0+16];1356; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_reassoc_nonpow2_param_0];1357; CHECK-SM100-NEXT: min.NaN.f32 %r8, %r4, %r5, %r6;1358; CHECK-SM100-NEXT: min.NaN.f32 %r9, %r1, %r2, %r3;1359; CHECK-SM100-NEXT: min.NaN.f32 %r10, %r9, %r8, %r7;1360; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r10;1361; CHECK-SM100-NEXT: ret;1362 %res = call reassoc float @llvm.vector.reduce.fminimum(<7 x float> %in)1363 ret float %res1364}1365 1366define i16 @reduce_add_i16(<8 x i16> %in) {1367; CHECK-SM80-LABEL: reduce_add_i16(1368; CHECK-SM80: {1369; CHECK-SM80-NEXT: .reg .b16 %rs<16>;1370; CHECK-SM80-NEXT: .reg .b32 %r<6>;1371; CHECK-SM80-EMPTY:1372; CHECK-SM80-NEXT: // %bb.0:1373; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_add_i16_param_0];1374; CHECK-SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;1375; CHECK-SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;1376; CHECK-SM80-NEXT: add.s16 %rs5, %rs4, %rs2;1377; CHECK-SM80-NEXT: mov.b32 {%rs6, %rs7}, %r3;1378; CHECK-SM80-NEXT: mov.b32 {%rs8, %rs9}, %r1;1379; CHECK-SM80-NEXT: add.s16 %rs10, %rs9, %rs7;1380; CHECK-SM80-NEXT: add.s16 %rs11, %rs10, %rs5;1381; CHECK-SM80-NEXT: add.s16 %rs12, %rs3, %rs1;1382; CHECK-SM80-NEXT: add.s16 %rs13, %rs8, %rs6;1383; CHECK-SM80-NEXT: add.s16 %rs14, %rs13, %rs12;1384; CHECK-SM80-NEXT: add.s16 %rs15, %rs14, %rs11;1385; CHECK-SM80-NEXT: cvt.u32.u16 %r5, %rs15;1386; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r5;1387; CHECK-SM80-NEXT: ret;1388;1389; CHECK-SM100-LABEL: reduce_add_i16(1390; CHECK-SM100: {1391; CHECK-SM100-NEXT: .reg .b16 %rs<4>;1392; CHECK-SM100-NEXT: .reg .b32 %r<9>;1393; CHECK-SM100-EMPTY:1394; CHECK-SM100-NEXT: // %bb.0:1395; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_add_i16_param_0];1396; CHECK-SM100-NEXT: add.s16x2 %r5, %r2, %r4;1397; CHECK-SM100-NEXT: add.s16x2 %r6, %r1, %r3;1398; CHECK-SM100-NEXT: add.s16x2 %r7, %r6, %r5;1399; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r7;1400; CHECK-SM100-NEXT: add.s16 %rs3, %rs1, %rs2;1401; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs3;1402; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1403; CHECK-SM100-NEXT: ret;1404 %res = call i16 @llvm.vector.reduce.add(<8 x i16> %in)1405 ret i16 %res1406}1407 1408define i16 @reduce_add_i16_nonpow2(<7 x i16> %in) {1409; CHECK-SM80-LABEL: reduce_add_i16_nonpow2(1410; CHECK-SM80: {1411; CHECK-SM80-NEXT: .reg .b16 %rs<14>;1412; CHECK-SM80-NEXT: .reg .b32 %r<2>;1413; CHECK-SM80-EMPTY:1414; CHECK-SM80-NEXT: // %bb.0:1415; CHECK-SM80-NEXT: ld.param.b16 %rs7, [reduce_add_i16_nonpow2_param_0+12];1416; CHECK-SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_add_i16_nonpow2_param_0+8];1417; CHECK-SM80-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_add_i16_nonpow2_param_0];1418; CHECK-SM80-NEXT: add.s16 %rs8, %rs3, %rs7;1419; CHECK-SM80-NEXT: add.s16 %rs9, %rs1, %rs5;1420; CHECK-SM80-NEXT: add.s16 %rs10, %rs9, %rs8;1421; CHECK-SM80-NEXT: add.s16 %rs11, %rs2, %rs6;1422; CHECK-SM80-NEXT: add.s16 %rs12, %rs11, %rs4;1423; CHECK-SM80-NEXT: add.s16 %rs13, %rs10, %rs12;1424; CHECK-SM80-NEXT: cvt.u32.u16 %r1, %rs13;1425; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r1;1426; CHECK-SM80-NEXT: ret;1427;1428; CHECK-SM100-LABEL: reduce_add_i16_nonpow2(1429; CHECK-SM100: {1430; CHECK-SM100-NEXT: .reg .b16 %rs<12>;1431; CHECK-SM100-NEXT: .reg .b32 %r<9>;1432; CHECK-SM100-EMPTY:1433; CHECK-SM100-NEXT: // %bb.0:1434; CHECK-SM100-NEXT: ld.param.b32 %r1, [reduce_add_i16_nonpow2_param_0+8];1435; CHECK-SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;1436; CHECK-SM100-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_add_i16_nonpow2_param_0];1437; CHECK-SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;1438; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r2;1439; CHECK-SM100-NEXT: ld.param.b16 %rs7, [reduce_add_i16_nonpow2_param_0+12];1440; CHECK-SM100-NEXT: add.s16x2 %r4, %r2, %r1;1441; CHECK-SM100-NEXT: mov.b16 %rs8, 0;1442; CHECK-SM100-NEXT: mov.b32 %r5, {%rs7, %rs8};1443; CHECK-SM100-NEXT: add.s16x2 %r6, %r3, %r5;1444; CHECK-SM100-NEXT: add.s16x2 %r7, %r4, %r6;1445; CHECK-SM100-NEXT: mov.b32 {%rs9, %rs10}, %r7;1446; CHECK-SM100-NEXT: add.s16 %rs11, %rs9, %rs10;1447; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs11;1448; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1449; CHECK-SM100-NEXT: ret;1450 %res = call i16 @llvm.vector.reduce.add(<7 x i16> %in)1451 ret i16 %res1452}1453 1454define i32 @reduce_add_i32(<8 x i32> %in) {1455; CHECK-SM80-LABEL: reduce_add_i32(1456; CHECK-SM80: {1457; CHECK-SM80-NEXT: .reg .b32 %r<16>;1458; CHECK-SM80-EMPTY:1459; CHECK-SM80-NEXT: // %bb.0:1460; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_add_i32_param_0+16];1461; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_add_i32_param_0];1462; CHECK-SM80-NEXT: add.s32 %r9, %r4, %r8;1463; CHECK-SM80-NEXT: add.s32 %r10, %r2, %r6;1464; CHECK-SM80-NEXT: add.s32 %r11, %r10, %r9;1465; CHECK-SM80-NEXT: add.s32 %r12, %r3, %r7;1466; CHECK-SM80-NEXT: add.s32 %r13, %r1, %r5;1467; CHECK-SM80-NEXT: add.s32 %r14, %r13, %r12;1468; CHECK-SM80-NEXT: add.s32 %r15, %r14, %r11;1469; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1470; CHECK-SM80-NEXT: ret;1471;1472; CHECK-SM100-LABEL: reduce_add_i32(1473; CHECK-SM100: {1474; CHECK-SM100-NEXT: .reg .b32 %r<16>;1475; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1476; CHECK-SM100-EMPTY:1477; CHECK-SM100-NEXT: // %bb.0:1478; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_add_i32_param_0+16];1479; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_add_i32_param_0];1480; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1481; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;1482; CHECK-SM100-NEXT: add.s32 %r5, %r4, %r2;1483; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;1484; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;1485; CHECK-SM100-NEXT: add.s32 %r10, %r9, %r7;1486; CHECK-SM100-NEXT: add.s32 %r11, %r10, %r5;1487; CHECK-SM100-NEXT: add.s32 %r12, %r3, %r1;1488; CHECK-SM100-NEXT: add.s32 %r13, %r8, %r6;1489; CHECK-SM100-NEXT: add.s32 %r14, %r13, %r12;1490; CHECK-SM100-NEXT: add.s32 %r15, %r14, %r11;1491; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;1492; CHECK-SM100-NEXT: ret;1493 %res = call i32 @llvm.vector.reduce.add(<8 x i32> %in)1494 ret i32 %res1495}1496 1497define i32 @reduce_add_i32_nonpow2(<7 x i32> %in) {1498; CHECK-LABEL: reduce_add_i32_nonpow2(1499; CHECK: {1500; CHECK-NEXT: .reg .b32 %r<14>;1501; CHECK-EMPTY:1502; CHECK-NEXT: // %bb.0:1503; CHECK-NEXT: ld.param.b32 %r7, [reduce_add_i32_nonpow2_param_0+24];1504; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_add_i32_nonpow2_param_0+16];1505; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_add_i32_nonpow2_param_0];1506; CHECK-NEXT: add.s32 %r8, %r3, %r7;1507; CHECK-NEXT: add.s32 %r9, %r1, %r5;1508; CHECK-NEXT: add.s32 %r10, %r9, %r8;1509; CHECK-NEXT: add.s32 %r11, %r2, %r6;1510; CHECK-NEXT: add.s32 %r12, %r11, %r4;1511; CHECK-NEXT: add.s32 %r13, %r10, %r12;1512; CHECK-NEXT: st.param.b32 [func_retval0], %r13;1513; CHECK-NEXT: ret;1514 %res = call i32 @llvm.vector.reduce.add(<7 x i32> %in)1515 ret i32 %res1516}1517 1518define i16 @reduce_mul_i16(<8 x i16> %in) {1519; CHECK-LABEL: reduce_mul_i16(1520; CHECK: {1521; CHECK-NEXT: .reg .b16 %rs<16>;1522; CHECK-NEXT: .reg .b32 %r<6>;1523; CHECK-EMPTY:1524; CHECK-NEXT: // %bb.0:1525; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_mul_i16_param_0];1526; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;1527; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r2;1528; CHECK-NEXT: mul.lo.s16 %rs5, %rs4, %rs2;1529; CHECK-NEXT: mov.b32 {%rs6, %rs7}, %r3;1530; CHECK-NEXT: mov.b32 {%rs8, %rs9}, %r1;1531; CHECK-NEXT: mul.lo.s16 %rs10, %rs9, %rs7;1532; CHECK-NEXT: mul.lo.s16 %rs11, %rs10, %rs5;1533; CHECK-NEXT: mul.lo.s16 %rs12, %rs3, %rs1;1534; CHECK-NEXT: mul.lo.s16 %rs13, %rs8, %rs6;1535; CHECK-NEXT: mul.lo.s16 %rs14, %rs13, %rs12;1536; CHECK-NEXT: mul.lo.s16 %rs15, %rs14, %rs11;1537; CHECK-NEXT: cvt.u32.u16 %r5, %rs15;1538; CHECK-NEXT: st.param.b32 [func_retval0], %r5;1539; CHECK-NEXT: ret;1540 %res = call i16 @llvm.vector.reduce.mul(<8 x i16> %in)1541 ret i16 %res1542}1543 1544define i16 @reduce_mul_i16_nonpow2(<7 x i16> %in) {1545; CHECK-LABEL: reduce_mul_i16_nonpow2(1546; CHECK: {1547; CHECK-NEXT: .reg .b16 %rs<14>;1548; CHECK-NEXT: .reg .b32 %r<2>;1549; CHECK-EMPTY:1550; CHECK-NEXT: // %bb.0:1551; CHECK-NEXT: ld.param.b16 %rs7, [reduce_mul_i16_nonpow2_param_0+12];1552; CHECK-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_mul_i16_nonpow2_param_0+8];1553; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_mul_i16_nonpow2_param_0];1554; CHECK-NEXT: mul.lo.s16 %rs8, %rs3, %rs7;1555; CHECK-NEXT: mul.lo.s16 %rs9, %rs1, %rs5;1556; CHECK-NEXT: mul.lo.s16 %rs10, %rs9, %rs8;1557; CHECK-NEXT: mul.lo.s16 %rs11, %rs2, %rs6;1558; CHECK-NEXT: mul.lo.s16 %rs12, %rs4, %rs11;1559; CHECK-NEXT: mul.lo.s16 %rs13, %rs10, %rs12;1560; CHECK-NEXT: cvt.u32.u16 %r1, %rs13;1561; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1562; CHECK-NEXT: ret;1563 %res = call i16 @llvm.vector.reduce.mul(<7 x i16> %in)1564 ret i16 %res1565}1566 1567define i32 @reduce_mul_i32(<8 x i32> %in) {1568; CHECK-SM80-LABEL: reduce_mul_i32(1569; CHECK-SM80: {1570; CHECK-SM80-NEXT: .reg .b32 %r<16>;1571; CHECK-SM80-EMPTY:1572; CHECK-SM80-NEXT: // %bb.0:1573; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_mul_i32_param_0+16];1574; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_mul_i32_param_0];1575; CHECK-SM80-NEXT: mul.lo.s32 %r9, %r4, %r8;1576; CHECK-SM80-NEXT: mul.lo.s32 %r10, %r2, %r6;1577; CHECK-SM80-NEXT: mul.lo.s32 %r11, %r10, %r9;1578; CHECK-SM80-NEXT: mul.lo.s32 %r12, %r3, %r7;1579; CHECK-SM80-NEXT: mul.lo.s32 %r13, %r1, %r5;1580; CHECK-SM80-NEXT: mul.lo.s32 %r14, %r13, %r12;1581; CHECK-SM80-NEXT: mul.lo.s32 %r15, %r14, %r11;1582; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1583; CHECK-SM80-NEXT: ret;1584;1585; CHECK-SM100-LABEL: reduce_mul_i32(1586; CHECK-SM100: {1587; CHECK-SM100-NEXT: .reg .b32 %r<16>;1588; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1589; CHECK-SM100-EMPTY:1590; CHECK-SM100-NEXT: // %bb.0:1591; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_mul_i32_param_0+16];1592; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_mul_i32_param_0];1593; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1594; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;1595; CHECK-SM100-NEXT: mul.lo.s32 %r5, %r4, %r2;1596; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;1597; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;1598; CHECK-SM100-NEXT: mul.lo.s32 %r10, %r9, %r7;1599; CHECK-SM100-NEXT: mul.lo.s32 %r11, %r10, %r5;1600; CHECK-SM100-NEXT: mul.lo.s32 %r12, %r3, %r1;1601; CHECK-SM100-NEXT: mul.lo.s32 %r13, %r8, %r6;1602; CHECK-SM100-NEXT: mul.lo.s32 %r14, %r13, %r12;1603; CHECK-SM100-NEXT: mul.lo.s32 %r15, %r14, %r11;1604; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;1605; CHECK-SM100-NEXT: ret;1606 %res = call i32 @llvm.vector.reduce.mul(<8 x i32> %in)1607 ret i32 %res1608}1609 1610define i32 @reduce_mul_i32_nonpow2(<7 x i32> %in) {1611; CHECK-LABEL: reduce_mul_i32_nonpow2(1612; CHECK: {1613; CHECK-NEXT: .reg .b32 %r<14>;1614; CHECK-EMPTY:1615; CHECK-NEXT: // %bb.0:1616; CHECK-NEXT: ld.param.b32 %r7, [reduce_mul_i32_nonpow2_param_0+24];1617; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_mul_i32_nonpow2_param_0+16];1618; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_mul_i32_nonpow2_param_0];1619; CHECK-NEXT: mul.lo.s32 %r8, %r3, %r7;1620; CHECK-NEXT: mul.lo.s32 %r9, %r1, %r5;1621; CHECK-NEXT: mul.lo.s32 %r10, %r9, %r8;1622; CHECK-NEXT: mul.lo.s32 %r11, %r2, %r6;1623; CHECK-NEXT: mul.lo.s32 %r12, %r4, %r11;1624; CHECK-NEXT: mul.lo.s32 %r13, %r10, %r12;1625; CHECK-NEXT: st.param.b32 [func_retval0], %r13;1626; CHECK-NEXT: ret;1627 %res = call i32 @llvm.vector.reduce.mul(<7 x i32> %in)1628 ret i32 %res1629}1630 1631define i16 @reduce_umax_i16(<8 x i16> %in) {1632; CHECK-SM80-LABEL: reduce_umax_i16(1633; CHECK-SM80: {1634; CHECK-SM80-NEXT: .reg .b16 %rs<16>;1635; CHECK-SM80-NEXT: .reg .b32 %r<6>;1636; CHECK-SM80-EMPTY:1637; CHECK-SM80-NEXT: // %bb.0:1638; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umax_i16_param_0];1639; CHECK-SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;1640; CHECK-SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;1641; CHECK-SM80-NEXT: max.u16 %rs5, %rs4, %rs2;1642; CHECK-SM80-NEXT: mov.b32 {%rs6, %rs7}, %r3;1643; CHECK-SM80-NEXT: mov.b32 {%rs8, %rs9}, %r1;1644; CHECK-SM80-NEXT: max.u16 %rs10, %rs9, %rs7;1645; CHECK-SM80-NEXT: max.u16 %rs11, %rs10, %rs5;1646; CHECK-SM80-NEXT: max.u16 %rs12, %rs3, %rs1;1647; CHECK-SM80-NEXT: max.u16 %rs13, %rs8, %rs6;1648; CHECK-SM80-NEXT: max.u16 %rs14, %rs13, %rs12;1649; CHECK-SM80-NEXT: max.u16 %rs15, %rs14, %rs11;1650; CHECK-SM80-NEXT: cvt.u32.u16 %r5, %rs15;1651; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r5;1652; CHECK-SM80-NEXT: ret;1653;1654; CHECK-SM100-LABEL: reduce_umax_i16(1655; CHECK-SM100: {1656; CHECK-SM100-NEXT: .reg .b16 %rs<4>;1657; CHECK-SM100-NEXT: .reg .b32 %r<9>;1658; CHECK-SM100-EMPTY:1659; CHECK-SM100-NEXT: // %bb.0:1660; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umax_i16_param_0];1661; CHECK-SM100-NEXT: max.u16x2 %r5, %r2, %r4;1662; CHECK-SM100-NEXT: max.u16x2 %r6, %r1, %r3;1663; CHECK-SM100-NEXT: max.u16x2 %r7, %r6, %r5;1664; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r7;1665; CHECK-SM100-NEXT: max.u16 %rs3, %rs1, %rs2;1666; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs3;1667; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1668; CHECK-SM100-NEXT: ret;1669 %res = call i16 @llvm.vector.reduce.umax(<8 x i16> %in)1670 ret i16 %res1671}1672 1673define i16 @reduce_umax_i16_nonpow2(<7 x i16> %in) {1674; CHECK-SM80-LABEL: reduce_umax_i16_nonpow2(1675; CHECK-SM80: {1676; CHECK-SM80-NEXT: .reg .b16 %rs<14>;1677; CHECK-SM80-NEXT: .reg .b32 %r<2>;1678; CHECK-SM80-EMPTY:1679; CHECK-SM80-NEXT: // %bb.0:1680; CHECK-SM80-NEXT: ld.param.b16 %rs7, [reduce_umax_i16_nonpow2_param_0+12];1681; CHECK-SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_umax_i16_nonpow2_param_0+8];1682; CHECK-SM80-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_umax_i16_nonpow2_param_0];1683; CHECK-SM80-NEXT: max.u16 %rs8, %rs3, %rs7;1684; CHECK-SM80-NEXT: max.u16 %rs9, %rs1, %rs5;1685; CHECK-SM80-NEXT: max.u16 %rs10, %rs9, %rs8;1686; CHECK-SM80-NEXT: max.u16 %rs11, %rs2, %rs6;1687; CHECK-SM80-NEXT: max.u16 %rs12, %rs4, %rs11;1688; CHECK-SM80-NEXT: max.u16 %rs13, %rs10, %rs12;1689; CHECK-SM80-NEXT: cvt.u32.u16 %r1, %rs13;1690; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r1;1691; CHECK-SM80-NEXT: ret;1692;1693; CHECK-SM100-LABEL: reduce_umax_i16_nonpow2(1694; CHECK-SM100: {1695; CHECK-SM100-NEXT: .reg .b16 %rs<12>;1696; CHECK-SM100-NEXT: .reg .b32 %r<9>;1697; CHECK-SM100-EMPTY:1698; CHECK-SM100-NEXT: // %bb.0:1699; CHECK-SM100-NEXT: ld.param.b32 %r1, [reduce_umax_i16_nonpow2_param_0+8];1700; CHECK-SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;1701; CHECK-SM100-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_umax_i16_nonpow2_param_0];1702; CHECK-SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;1703; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r2;1704; CHECK-SM100-NEXT: ld.param.b16 %rs7, [reduce_umax_i16_nonpow2_param_0+12];1705; CHECK-SM100-NEXT: max.u16x2 %r4, %r2, %r1;1706; CHECK-SM100-NEXT: mov.b16 %rs8, 0;1707; CHECK-SM100-NEXT: mov.b32 %r5, {%rs7, %rs8};1708; CHECK-SM100-NEXT: max.u16x2 %r6, %r3, %r5;1709; CHECK-SM100-NEXT: max.u16x2 %r7, %r4, %r6;1710; CHECK-SM100-NEXT: mov.b32 {%rs9, %rs10}, %r7;1711; CHECK-SM100-NEXT: max.u16 %rs11, %rs9, %rs10;1712; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs11;1713; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1714; CHECK-SM100-NEXT: ret;1715 %res = call i16 @llvm.vector.reduce.umax(<7 x i16> %in)1716 ret i16 %res1717}1718 1719define i32 @reduce_umax_i32(<8 x i32> %in) {1720; CHECK-SM80-LABEL: reduce_umax_i32(1721; CHECK-SM80: {1722; CHECK-SM80-NEXT: .reg .b32 %r<16>;1723; CHECK-SM80-EMPTY:1724; CHECK-SM80-NEXT: // %bb.0:1725; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_umax_i32_param_0+16];1726; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umax_i32_param_0];1727; CHECK-SM80-NEXT: max.u32 %r9, %r4, %r8;1728; CHECK-SM80-NEXT: max.u32 %r10, %r2, %r6;1729; CHECK-SM80-NEXT: max.u32 %r11, %r10, %r9;1730; CHECK-SM80-NEXT: max.u32 %r12, %r3, %r7;1731; CHECK-SM80-NEXT: max.u32 %r13, %r1, %r5;1732; CHECK-SM80-NEXT: max.u32 %r14, %r13, %r12;1733; CHECK-SM80-NEXT: max.u32 %r15, %r14, %r11;1734; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1735; CHECK-SM80-NEXT: ret;1736;1737; CHECK-SM100-LABEL: reduce_umax_i32(1738; CHECK-SM100: {1739; CHECK-SM100-NEXT: .reg .b32 %r<16>;1740; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1741; CHECK-SM100-EMPTY:1742; CHECK-SM100-NEXT: // %bb.0:1743; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_umax_i32_param_0+16];1744; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_umax_i32_param_0];1745; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1746; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;1747; CHECK-SM100-NEXT: max.u32 %r5, %r4, %r2;1748; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;1749; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;1750; CHECK-SM100-NEXT: max.u32 %r10, %r9, %r7;1751; CHECK-SM100-NEXT: max.u32 %r11, %r10, %r5;1752; CHECK-SM100-NEXT: max.u32 %r12, %r3, %r1;1753; CHECK-SM100-NEXT: max.u32 %r13, %r8, %r6;1754; CHECK-SM100-NEXT: max.u32 %r14, %r13, %r12;1755; CHECK-SM100-NEXT: max.u32 %r15, %r14, %r11;1756; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;1757; CHECK-SM100-NEXT: ret;1758 %res = call i32 @llvm.vector.reduce.umax(<8 x i32> %in)1759 ret i32 %res1760}1761 1762define i32 @reduce_umax_i32_nonpow2(<7 x i32> %in) {1763; CHECK-LABEL: reduce_umax_i32_nonpow2(1764; CHECK: {1765; CHECK-NEXT: .reg .b32 %r<14>;1766; CHECK-EMPTY:1767; CHECK-NEXT: // %bb.0:1768; CHECK-NEXT: ld.param.b32 %r7, [reduce_umax_i32_nonpow2_param_0+24];1769; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_umax_i32_nonpow2_param_0+16];1770; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umax_i32_nonpow2_param_0];1771; CHECK-NEXT: max.u32 %r8, %r3, %r7;1772; CHECK-NEXT: max.u32 %r9, %r1, %r5;1773; CHECK-NEXT: max.u32 %r10, %r9, %r8;1774; CHECK-NEXT: max.u32 %r11, %r2, %r6;1775; CHECK-NEXT: max.u32 %r12, %r4, %r11;1776; CHECK-NEXT: max.u32 %r13, %r10, %r12;1777; CHECK-NEXT: st.param.b32 [func_retval0], %r13;1778; CHECK-NEXT: ret;1779 %res = call i32 @llvm.vector.reduce.umax(<7 x i32> %in)1780 ret i32 %res1781}1782 1783define i16 @reduce_umin_i16(<8 x i16> %in) {1784; CHECK-SM80-LABEL: reduce_umin_i16(1785; CHECK-SM80: {1786; CHECK-SM80-NEXT: .reg .b16 %rs<16>;1787; CHECK-SM80-NEXT: .reg .b32 %r<6>;1788; CHECK-SM80-EMPTY:1789; CHECK-SM80-NEXT: // %bb.0:1790; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umin_i16_param_0];1791; CHECK-SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;1792; CHECK-SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;1793; CHECK-SM80-NEXT: min.u16 %rs5, %rs4, %rs2;1794; CHECK-SM80-NEXT: mov.b32 {%rs6, %rs7}, %r3;1795; CHECK-SM80-NEXT: mov.b32 {%rs8, %rs9}, %r1;1796; CHECK-SM80-NEXT: min.u16 %rs10, %rs9, %rs7;1797; CHECK-SM80-NEXT: min.u16 %rs11, %rs10, %rs5;1798; CHECK-SM80-NEXT: min.u16 %rs12, %rs3, %rs1;1799; CHECK-SM80-NEXT: min.u16 %rs13, %rs8, %rs6;1800; CHECK-SM80-NEXT: min.u16 %rs14, %rs13, %rs12;1801; CHECK-SM80-NEXT: min.u16 %rs15, %rs14, %rs11;1802; CHECK-SM80-NEXT: cvt.u32.u16 %r5, %rs15;1803; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r5;1804; CHECK-SM80-NEXT: ret;1805;1806; CHECK-SM100-LABEL: reduce_umin_i16(1807; CHECK-SM100: {1808; CHECK-SM100-NEXT: .reg .b16 %rs<4>;1809; CHECK-SM100-NEXT: .reg .b32 %r<9>;1810; CHECK-SM100-EMPTY:1811; CHECK-SM100-NEXT: // %bb.0:1812; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umin_i16_param_0];1813; CHECK-SM100-NEXT: min.u16x2 %r5, %r2, %r4;1814; CHECK-SM100-NEXT: min.u16x2 %r6, %r1, %r3;1815; CHECK-SM100-NEXT: min.u16x2 %r7, %r6, %r5;1816; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r7;1817; CHECK-SM100-NEXT: min.u16 %rs3, %rs1, %rs2;1818; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs3;1819; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1820; CHECK-SM100-NEXT: ret;1821 %res = call i16 @llvm.vector.reduce.umin(<8 x i16> %in)1822 ret i16 %res1823}1824 1825define i16 @reduce_umin_i16_nonpow2(<7 x i16> %in) {1826; CHECK-SM80-LABEL: reduce_umin_i16_nonpow2(1827; CHECK-SM80: {1828; CHECK-SM80-NEXT: .reg .b16 %rs<14>;1829; CHECK-SM80-NEXT: .reg .b32 %r<2>;1830; CHECK-SM80-EMPTY:1831; CHECK-SM80-NEXT: // %bb.0:1832; CHECK-SM80-NEXT: ld.param.b16 %rs7, [reduce_umin_i16_nonpow2_param_0+12];1833; CHECK-SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_umin_i16_nonpow2_param_0+8];1834; CHECK-SM80-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_umin_i16_nonpow2_param_0];1835; CHECK-SM80-NEXT: min.u16 %rs8, %rs3, %rs7;1836; CHECK-SM80-NEXT: min.u16 %rs9, %rs1, %rs5;1837; CHECK-SM80-NEXT: min.u16 %rs10, %rs9, %rs8;1838; CHECK-SM80-NEXT: min.u16 %rs11, %rs2, %rs6;1839; CHECK-SM80-NEXT: min.u16 %rs12, %rs4, %rs11;1840; CHECK-SM80-NEXT: min.u16 %rs13, %rs10, %rs12;1841; CHECK-SM80-NEXT: cvt.u32.u16 %r1, %rs13;1842; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r1;1843; CHECK-SM80-NEXT: ret;1844;1845; CHECK-SM100-LABEL: reduce_umin_i16_nonpow2(1846; CHECK-SM100: {1847; CHECK-SM100-NEXT: .reg .b16 %rs<12>;1848; CHECK-SM100-NEXT: .reg .b32 %r<9>;1849; CHECK-SM100-EMPTY:1850; CHECK-SM100-NEXT: // %bb.0:1851; CHECK-SM100-NEXT: ld.param.b32 %r1, [reduce_umin_i16_nonpow2_param_0+8];1852; CHECK-SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;1853; CHECK-SM100-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_umin_i16_nonpow2_param_0];1854; CHECK-SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;1855; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r2;1856; CHECK-SM100-NEXT: ld.param.b16 %rs7, [reduce_umin_i16_nonpow2_param_0+12];1857; CHECK-SM100-NEXT: min.u16x2 %r4, %r2, %r1;1858; CHECK-SM100-NEXT: mov.b16 %rs8, -1;1859; CHECK-SM100-NEXT: mov.b32 %r5, {%rs7, %rs8};1860; CHECK-SM100-NEXT: min.u16x2 %r6, %r3, %r5;1861; CHECK-SM100-NEXT: min.u16x2 %r7, %r4, %r6;1862; CHECK-SM100-NEXT: mov.b32 {%rs9, %rs10}, %r7;1863; CHECK-SM100-NEXT: min.u16 %rs11, %rs9, %rs10;1864; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs11;1865; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1866; CHECK-SM100-NEXT: ret;1867 %res = call i16 @llvm.vector.reduce.umin(<7 x i16> %in)1868 ret i16 %res1869}1870 1871define i32 @reduce_umin_i32(<8 x i32> %in) {1872; CHECK-SM80-LABEL: reduce_umin_i32(1873; CHECK-SM80: {1874; CHECK-SM80-NEXT: .reg .b32 %r<16>;1875; CHECK-SM80-EMPTY:1876; CHECK-SM80-NEXT: // %bb.0:1877; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_umin_i32_param_0+16];1878; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umin_i32_param_0];1879; CHECK-SM80-NEXT: min.u32 %r9, %r4, %r8;1880; CHECK-SM80-NEXT: min.u32 %r10, %r2, %r6;1881; CHECK-SM80-NEXT: min.u32 %r11, %r10, %r9;1882; CHECK-SM80-NEXT: min.u32 %r12, %r3, %r7;1883; CHECK-SM80-NEXT: min.u32 %r13, %r1, %r5;1884; CHECK-SM80-NEXT: min.u32 %r14, %r13, %r12;1885; CHECK-SM80-NEXT: min.u32 %r15, %r14, %r11;1886; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;1887; CHECK-SM80-NEXT: ret;1888;1889; CHECK-SM100-LABEL: reduce_umin_i32(1890; CHECK-SM100: {1891; CHECK-SM100-NEXT: .reg .b32 %r<16>;1892; CHECK-SM100-NEXT: .reg .b64 %rd<5>;1893; CHECK-SM100-EMPTY:1894; CHECK-SM100-NEXT: // %bb.0:1895; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_umin_i32_param_0+16];1896; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_umin_i32_param_0];1897; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;1898; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;1899; CHECK-SM100-NEXT: min.u32 %r5, %r4, %r2;1900; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;1901; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;1902; CHECK-SM100-NEXT: min.u32 %r10, %r9, %r7;1903; CHECK-SM100-NEXT: min.u32 %r11, %r10, %r5;1904; CHECK-SM100-NEXT: min.u32 %r12, %r3, %r1;1905; CHECK-SM100-NEXT: min.u32 %r13, %r8, %r6;1906; CHECK-SM100-NEXT: min.u32 %r14, %r13, %r12;1907; CHECK-SM100-NEXT: min.u32 %r15, %r14, %r11;1908; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;1909; CHECK-SM100-NEXT: ret;1910 %res = call i32 @llvm.vector.reduce.umin(<8 x i32> %in)1911 ret i32 %res1912}1913 1914define i32 @reduce_umin_i32_nonpow2(<7 x i32> %in) {1915; CHECK-LABEL: reduce_umin_i32_nonpow2(1916; CHECK: {1917; CHECK-NEXT: .reg .b32 %r<14>;1918; CHECK-EMPTY:1919; CHECK-NEXT: // %bb.0:1920; CHECK-NEXT: ld.param.b32 %r7, [reduce_umin_i32_nonpow2_param_0+24];1921; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_umin_i32_nonpow2_param_0+16];1922; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_umin_i32_nonpow2_param_0];1923; CHECK-NEXT: min.u32 %r8, %r3, %r7;1924; CHECK-NEXT: min.u32 %r9, %r1, %r5;1925; CHECK-NEXT: min.u32 %r10, %r9, %r8;1926; CHECK-NEXT: min.u32 %r11, %r2, %r6;1927; CHECK-NEXT: min.u32 %r12, %r4, %r11;1928; CHECK-NEXT: min.u32 %r13, %r10, %r12;1929; CHECK-NEXT: st.param.b32 [func_retval0], %r13;1930; CHECK-NEXT: ret;1931 %res = call i32 @llvm.vector.reduce.umin(<7 x i32> %in)1932 ret i32 %res1933}1934 1935define i16 @reduce_smax_i16(<8 x i16> %in) {1936; CHECK-SM80-LABEL: reduce_smax_i16(1937; CHECK-SM80: {1938; CHECK-SM80-NEXT: .reg .b16 %rs<16>;1939; CHECK-SM80-NEXT: .reg .b32 %r<6>;1940; CHECK-SM80-EMPTY:1941; CHECK-SM80-NEXT: // %bb.0:1942; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smax_i16_param_0];1943; CHECK-SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;1944; CHECK-SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;1945; CHECK-SM80-NEXT: max.s16 %rs5, %rs4, %rs2;1946; CHECK-SM80-NEXT: mov.b32 {%rs6, %rs7}, %r3;1947; CHECK-SM80-NEXT: mov.b32 {%rs8, %rs9}, %r1;1948; CHECK-SM80-NEXT: max.s16 %rs10, %rs9, %rs7;1949; CHECK-SM80-NEXT: max.s16 %rs11, %rs10, %rs5;1950; CHECK-SM80-NEXT: max.s16 %rs12, %rs3, %rs1;1951; CHECK-SM80-NEXT: max.s16 %rs13, %rs8, %rs6;1952; CHECK-SM80-NEXT: max.s16 %rs14, %rs13, %rs12;1953; CHECK-SM80-NEXT: max.s16 %rs15, %rs14, %rs11;1954; CHECK-SM80-NEXT: cvt.u32.u16 %r5, %rs15;1955; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r5;1956; CHECK-SM80-NEXT: ret;1957;1958; CHECK-SM100-LABEL: reduce_smax_i16(1959; CHECK-SM100: {1960; CHECK-SM100-NEXT: .reg .b16 %rs<4>;1961; CHECK-SM100-NEXT: .reg .b32 %r<9>;1962; CHECK-SM100-EMPTY:1963; CHECK-SM100-NEXT: // %bb.0:1964; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smax_i16_param_0];1965; CHECK-SM100-NEXT: max.s16x2 %r5, %r2, %r4;1966; CHECK-SM100-NEXT: max.s16x2 %r6, %r1, %r3;1967; CHECK-SM100-NEXT: max.s16x2 %r7, %r6, %r5;1968; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r7;1969; CHECK-SM100-NEXT: max.s16 %rs3, %rs1, %rs2;1970; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs3;1971; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;1972; CHECK-SM100-NEXT: ret;1973 %res = call i16 @llvm.vector.reduce.smax(<8 x i16> %in)1974 ret i16 %res1975}1976 1977define i16 @reduce_smax_i16_nonpow2(<7 x i16> %in) {1978; CHECK-SM80-LABEL: reduce_smax_i16_nonpow2(1979; CHECK-SM80: {1980; CHECK-SM80-NEXT: .reg .b16 %rs<14>;1981; CHECK-SM80-NEXT: .reg .b32 %r<2>;1982; CHECK-SM80-EMPTY:1983; CHECK-SM80-NEXT: // %bb.0:1984; CHECK-SM80-NEXT: ld.param.b16 %rs7, [reduce_smax_i16_nonpow2_param_0+12];1985; CHECK-SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_smax_i16_nonpow2_param_0+8];1986; CHECK-SM80-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_smax_i16_nonpow2_param_0];1987; CHECK-SM80-NEXT: max.s16 %rs8, %rs3, %rs7;1988; CHECK-SM80-NEXT: max.s16 %rs9, %rs1, %rs5;1989; CHECK-SM80-NEXT: max.s16 %rs10, %rs9, %rs8;1990; CHECK-SM80-NEXT: max.s16 %rs11, %rs2, %rs6;1991; CHECK-SM80-NEXT: max.s16 %rs12, %rs4, %rs11;1992; CHECK-SM80-NEXT: max.s16 %rs13, %rs10, %rs12;1993; CHECK-SM80-NEXT: cvt.u32.u16 %r1, %rs13;1994; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r1;1995; CHECK-SM80-NEXT: ret;1996;1997; CHECK-SM100-LABEL: reduce_smax_i16_nonpow2(1998; CHECK-SM100: {1999; CHECK-SM100-NEXT: .reg .b16 %rs<12>;2000; CHECK-SM100-NEXT: .reg .b32 %r<9>;2001; CHECK-SM100-EMPTY:2002; CHECK-SM100-NEXT: // %bb.0:2003; CHECK-SM100-NEXT: ld.param.b32 %r1, [reduce_smax_i16_nonpow2_param_0+8];2004; CHECK-SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;2005; CHECK-SM100-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_smax_i16_nonpow2_param_0];2006; CHECK-SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;2007; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r2;2008; CHECK-SM100-NEXT: ld.param.b16 %rs7, [reduce_smax_i16_nonpow2_param_0+12];2009; CHECK-SM100-NEXT: max.s16x2 %r4, %r2, %r1;2010; CHECK-SM100-NEXT: mov.b16 %rs8, -32768;2011; CHECK-SM100-NEXT: mov.b32 %r5, {%rs7, %rs8};2012; CHECK-SM100-NEXT: max.s16x2 %r6, %r3, %r5;2013; CHECK-SM100-NEXT: max.s16x2 %r7, %r4, %r6;2014; CHECK-SM100-NEXT: mov.b32 {%rs9, %rs10}, %r7;2015; CHECK-SM100-NEXT: max.s16 %rs11, %rs9, %rs10;2016; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs11;2017; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;2018; CHECK-SM100-NEXT: ret;2019 %res = call i16 @llvm.vector.reduce.smax(<7 x i16> %in)2020 ret i16 %res2021}2022 2023define i32 @reduce_smax_i32(<8 x i32> %in) {2024; CHECK-SM80-LABEL: reduce_smax_i32(2025; CHECK-SM80: {2026; CHECK-SM80-NEXT: .reg .b32 %r<16>;2027; CHECK-SM80-EMPTY:2028; CHECK-SM80-NEXT: // %bb.0:2029; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_smax_i32_param_0+16];2030; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smax_i32_param_0];2031; CHECK-SM80-NEXT: max.s32 %r9, %r4, %r8;2032; CHECK-SM80-NEXT: max.s32 %r10, %r2, %r6;2033; CHECK-SM80-NEXT: max.s32 %r11, %r10, %r9;2034; CHECK-SM80-NEXT: max.s32 %r12, %r3, %r7;2035; CHECK-SM80-NEXT: max.s32 %r13, %r1, %r5;2036; CHECK-SM80-NEXT: max.s32 %r14, %r13, %r12;2037; CHECK-SM80-NEXT: max.s32 %r15, %r14, %r11;2038; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;2039; CHECK-SM80-NEXT: ret;2040;2041; CHECK-SM100-LABEL: reduce_smax_i32(2042; CHECK-SM100: {2043; CHECK-SM100-NEXT: .reg .b32 %r<16>;2044; CHECK-SM100-NEXT: .reg .b64 %rd<5>;2045; CHECK-SM100-EMPTY:2046; CHECK-SM100-NEXT: // %bb.0:2047; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_smax_i32_param_0+16];2048; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_smax_i32_param_0];2049; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;2050; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;2051; CHECK-SM100-NEXT: max.s32 %r5, %r4, %r2;2052; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;2053; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;2054; CHECK-SM100-NEXT: max.s32 %r10, %r9, %r7;2055; CHECK-SM100-NEXT: max.s32 %r11, %r10, %r5;2056; CHECK-SM100-NEXT: max.s32 %r12, %r3, %r1;2057; CHECK-SM100-NEXT: max.s32 %r13, %r8, %r6;2058; CHECK-SM100-NEXT: max.s32 %r14, %r13, %r12;2059; CHECK-SM100-NEXT: max.s32 %r15, %r14, %r11;2060; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;2061; CHECK-SM100-NEXT: ret;2062 %res = call i32 @llvm.vector.reduce.smax(<8 x i32> %in)2063 ret i32 %res2064}2065 2066define i32 @reduce_smax_i32_nonpow2(<7 x i32> %in) {2067; CHECK-LABEL: reduce_smax_i32_nonpow2(2068; CHECK: {2069; CHECK-NEXT: .reg .b32 %r<14>;2070; CHECK-EMPTY:2071; CHECK-NEXT: // %bb.0:2072; CHECK-NEXT: ld.param.b32 %r7, [reduce_smax_i32_nonpow2_param_0+24];2073; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_smax_i32_nonpow2_param_0+16];2074; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smax_i32_nonpow2_param_0];2075; CHECK-NEXT: max.s32 %r8, %r3, %r7;2076; CHECK-NEXT: max.s32 %r9, %r1, %r5;2077; CHECK-NEXT: max.s32 %r10, %r9, %r8;2078; CHECK-NEXT: max.s32 %r11, %r2, %r6;2079; CHECK-NEXT: max.s32 %r12, %r4, %r11;2080; CHECK-NEXT: max.s32 %r13, %r10, %r12;2081; CHECK-NEXT: st.param.b32 [func_retval0], %r13;2082; CHECK-NEXT: ret;2083 %res = call i32 @llvm.vector.reduce.smax(<7 x i32> %in)2084 ret i32 %res2085}2086 2087define i16 @reduce_smin_i16(<8 x i16> %in) {2088; CHECK-SM80-LABEL: reduce_smin_i16(2089; CHECK-SM80: {2090; CHECK-SM80-NEXT: .reg .b16 %rs<16>;2091; CHECK-SM80-NEXT: .reg .b32 %r<6>;2092; CHECK-SM80-EMPTY:2093; CHECK-SM80-NEXT: // %bb.0:2094; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smin_i16_param_0];2095; CHECK-SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;2096; CHECK-SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;2097; CHECK-SM80-NEXT: min.s16 %rs5, %rs4, %rs2;2098; CHECK-SM80-NEXT: mov.b32 {%rs6, %rs7}, %r3;2099; CHECK-SM80-NEXT: mov.b32 {%rs8, %rs9}, %r1;2100; CHECK-SM80-NEXT: min.s16 %rs10, %rs9, %rs7;2101; CHECK-SM80-NEXT: min.s16 %rs11, %rs10, %rs5;2102; CHECK-SM80-NEXT: min.s16 %rs12, %rs3, %rs1;2103; CHECK-SM80-NEXT: min.s16 %rs13, %rs8, %rs6;2104; CHECK-SM80-NEXT: min.s16 %rs14, %rs13, %rs12;2105; CHECK-SM80-NEXT: min.s16 %rs15, %rs14, %rs11;2106; CHECK-SM80-NEXT: cvt.u32.u16 %r5, %rs15;2107; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r5;2108; CHECK-SM80-NEXT: ret;2109;2110; CHECK-SM100-LABEL: reduce_smin_i16(2111; CHECK-SM100: {2112; CHECK-SM100-NEXT: .reg .b16 %rs<4>;2113; CHECK-SM100-NEXT: .reg .b32 %r<9>;2114; CHECK-SM100-EMPTY:2115; CHECK-SM100-NEXT: // %bb.0:2116; CHECK-SM100-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smin_i16_param_0];2117; CHECK-SM100-NEXT: min.s16x2 %r5, %r2, %r4;2118; CHECK-SM100-NEXT: min.s16x2 %r6, %r1, %r3;2119; CHECK-SM100-NEXT: min.s16x2 %r7, %r6, %r5;2120; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r7;2121; CHECK-SM100-NEXT: min.s16 %rs3, %rs1, %rs2;2122; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs3;2123; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;2124; CHECK-SM100-NEXT: ret;2125 %res = call i16 @llvm.vector.reduce.smin(<8 x i16> %in)2126 ret i16 %res2127}2128 2129define i16 @reduce_smin_i16_nonpow2(<7 x i16> %in) {2130; CHECK-SM80-LABEL: reduce_smin_i16_nonpow2(2131; CHECK-SM80: {2132; CHECK-SM80-NEXT: .reg .b16 %rs<14>;2133; CHECK-SM80-NEXT: .reg .b32 %r<2>;2134; CHECK-SM80-EMPTY:2135; CHECK-SM80-NEXT: // %bb.0:2136; CHECK-SM80-NEXT: ld.param.b16 %rs7, [reduce_smin_i16_nonpow2_param_0+12];2137; CHECK-SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [reduce_smin_i16_nonpow2_param_0+8];2138; CHECK-SM80-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [reduce_smin_i16_nonpow2_param_0];2139; CHECK-SM80-NEXT: min.s16 %rs8, %rs3, %rs7;2140; CHECK-SM80-NEXT: min.s16 %rs9, %rs1, %rs5;2141; CHECK-SM80-NEXT: min.s16 %rs10, %rs9, %rs8;2142; CHECK-SM80-NEXT: min.s16 %rs11, %rs2, %rs6;2143; CHECK-SM80-NEXT: min.s16 %rs12, %rs4, %rs11;2144; CHECK-SM80-NEXT: min.s16 %rs13, %rs10, %rs12;2145; CHECK-SM80-NEXT: cvt.u32.u16 %r1, %rs13;2146; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r1;2147; CHECK-SM80-NEXT: ret;2148;2149; CHECK-SM100-LABEL: reduce_smin_i16_nonpow2(2150; CHECK-SM100: {2151; CHECK-SM100-NEXT: .reg .b16 %rs<12>;2152; CHECK-SM100-NEXT: .reg .b32 %r<9>;2153; CHECK-SM100-EMPTY:2154; CHECK-SM100-NEXT: // %bb.0:2155; CHECK-SM100-NEXT: ld.param.b32 %r1, [reduce_smin_i16_nonpow2_param_0+8];2156; CHECK-SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;2157; CHECK-SM100-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_smin_i16_nonpow2_param_0];2158; CHECK-SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;2159; CHECK-SM100-NEXT: mov.b32 {%rs1, %rs2}, %r2;2160; CHECK-SM100-NEXT: ld.param.b16 %rs7, [reduce_smin_i16_nonpow2_param_0+12];2161; CHECK-SM100-NEXT: min.s16x2 %r4, %r2, %r1;2162; CHECK-SM100-NEXT: mov.b16 %rs8, 32767;2163; CHECK-SM100-NEXT: mov.b32 %r5, {%rs7, %rs8};2164; CHECK-SM100-NEXT: min.s16x2 %r6, %r3, %r5;2165; CHECK-SM100-NEXT: min.s16x2 %r7, %r4, %r6;2166; CHECK-SM100-NEXT: mov.b32 {%rs9, %rs10}, %r7;2167; CHECK-SM100-NEXT: min.s16 %rs11, %rs9, %rs10;2168; CHECK-SM100-NEXT: cvt.u32.u16 %r8, %rs11;2169; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r8;2170; CHECK-SM100-NEXT: ret;2171 %res = call i16 @llvm.vector.reduce.smin(<7 x i16> %in)2172 ret i16 %res2173}2174 2175define i32 @reduce_smin_i32(<8 x i32> %in) {2176; CHECK-SM80-LABEL: reduce_smin_i32(2177; CHECK-SM80: {2178; CHECK-SM80-NEXT: .reg .b32 %r<16>;2179; CHECK-SM80-EMPTY:2180; CHECK-SM80-NEXT: // %bb.0:2181; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_smin_i32_param_0+16];2182; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smin_i32_param_0];2183; CHECK-SM80-NEXT: min.s32 %r9, %r4, %r8;2184; CHECK-SM80-NEXT: min.s32 %r10, %r2, %r6;2185; CHECK-SM80-NEXT: min.s32 %r11, %r10, %r9;2186; CHECK-SM80-NEXT: min.s32 %r12, %r3, %r7;2187; CHECK-SM80-NEXT: min.s32 %r13, %r1, %r5;2188; CHECK-SM80-NEXT: min.s32 %r14, %r13, %r12;2189; CHECK-SM80-NEXT: min.s32 %r15, %r14, %r11;2190; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;2191; CHECK-SM80-NEXT: ret;2192;2193; CHECK-SM100-LABEL: reduce_smin_i32(2194; CHECK-SM100: {2195; CHECK-SM100-NEXT: .reg .b32 %r<16>;2196; CHECK-SM100-NEXT: .reg .b64 %rd<5>;2197; CHECK-SM100-EMPTY:2198; CHECK-SM100-NEXT: // %bb.0:2199; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_smin_i32_param_0+16];2200; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_smin_i32_param_0];2201; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;2202; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;2203; CHECK-SM100-NEXT: min.s32 %r5, %r4, %r2;2204; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;2205; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;2206; CHECK-SM100-NEXT: min.s32 %r10, %r9, %r7;2207; CHECK-SM100-NEXT: min.s32 %r11, %r10, %r5;2208; CHECK-SM100-NEXT: min.s32 %r12, %r3, %r1;2209; CHECK-SM100-NEXT: min.s32 %r13, %r8, %r6;2210; CHECK-SM100-NEXT: min.s32 %r14, %r13, %r12;2211; CHECK-SM100-NEXT: min.s32 %r15, %r14, %r11;2212; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;2213; CHECK-SM100-NEXT: ret;2214 %res = call i32 @llvm.vector.reduce.smin(<8 x i32> %in)2215 ret i32 %res2216}2217 2218define i32 @reduce_smin_i32_nonpow2(<7 x i32> %in) {2219; CHECK-LABEL: reduce_smin_i32_nonpow2(2220; CHECK: {2221; CHECK-NEXT: .reg .b32 %r<14>;2222; CHECK-EMPTY:2223; CHECK-NEXT: // %bb.0:2224; CHECK-NEXT: ld.param.b32 %r7, [reduce_smin_i32_nonpow2_param_0+24];2225; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_smin_i32_nonpow2_param_0+16];2226; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_smin_i32_nonpow2_param_0];2227; CHECK-NEXT: min.s32 %r8, %r3, %r7;2228; CHECK-NEXT: min.s32 %r9, %r1, %r5;2229; CHECK-NEXT: min.s32 %r10, %r9, %r8;2230; CHECK-NEXT: min.s32 %r11, %r2, %r6;2231; CHECK-NEXT: min.s32 %r12, %r4, %r11;2232; CHECK-NEXT: min.s32 %r13, %r10, %r12;2233; CHECK-NEXT: st.param.b32 [func_retval0], %r13;2234; CHECK-NEXT: ret;2235 %res = call i32 @llvm.vector.reduce.smin(<7 x i32> %in)2236 ret i32 %res2237}2238 2239define i16 @reduce_and_i16(<8 x i16> %in) {2240; CHECK-LABEL: reduce_and_i16(2241; CHECK: {2242; CHECK-NEXT: .reg .b16 %rs<4>;2243; CHECK-NEXT: .reg .b32 %r<9>;2244; CHECK-EMPTY:2245; CHECK-NEXT: // %bb.0:2246; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_and_i16_param_0];2247; CHECK-NEXT: and.b32 %r5, %r2, %r4;2248; CHECK-NEXT: and.b32 %r6, %r1, %r3;2249; CHECK-NEXT: and.b32 %r7, %r6, %r5;2250; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;2251; CHECK-NEXT: and.b16 %rs3, %rs1, %rs2;2252; CHECK-NEXT: cvt.u32.u16 %r8, %rs3;2253; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2254; CHECK-NEXT: ret;2255 %res = call i16 @llvm.vector.reduce.and(<8 x i16> %in)2256 ret i16 %res2257}2258 2259define i16 @reduce_and_i16_nonpow2(<7 x i16> %in) {2260; CHECK-LABEL: reduce_and_i16_nonpow2(2261; CHECK: {2262; CHECK-NEXT: .reg .b16 %rs<12>;2263; CHECK-NEXT: .reg .b32 %r<9>;2264; CHECK-EMPTY:2265; CHECK-NEXT: // %bb.0:2266; CHECK-NEXT: ld.param.b32 %r1, [reduce_and_i16_nonpow2_param_0+8];2267; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;2268; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_and_i16_nonpow2_param_0];2269; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;2270; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;2271; CHECK-NEXT: ld.param.b16 %rs7, [reduce_and_i16_nonpow2_param_0+12];2272; CHECK-NEXT: and.b32 %r4, %r2, %r1;2273; CHECK-NEXT: mov.b16 %rs8, -1;2274; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};2275; CHECK-NEXT: and.b32 %r6, %r3, %r5;2276; CHECK-NEXT: and.b32 %r7, %r4, %r6;2277; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;2278; CHECK-NEXT: and.b16 %rs11, %rs9, %rs10;2279; CHECK-NEXT: cvt.u32.u16 %r8, %rs11;2280; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2281; CHECK-NEXT: ret;2282 %res = call i16 @llvm.vector.reduce.and(<7 x i16> %in)2283 ret i16 %res2284}2285 2286define i32 @reduce_and_i32(<8 x i32> %in) {2287; CHECK-SM80-LABEL: reduce_and_i32(2288; CHECK-SM80: {2289; CHECK-SM80-NEXT: .reg .b32 %r<16>;2290; CHECK-SM80-EMPTY:2291; CHECK-SM80-NEXT: // %bb.0:2292; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_and_i32_param_0+16];2293; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_and_i32_param_0];2294; CHECK-SM80-NEXT: and.b32 %r9, %r4, %r8;2295; CHECK-SM80-NEXT: and.b32 %r10, %r2, %r6;2296; CHECK-SM80-NEXT: and.b32 %r11, %r10, %r9;2297; CHECK-SM80-NEXT: and.b32 %r12, %r3, %r7;2298; CHECK-SM80-NEXT: and.b32 %r13, %r1, %r5;2299; CHECK-SM80-NEXT: and.b32 %r14, %r13, %r12;2300; CHECK-SM80-NEXT: and.b32 %r15, %r14, %r11;2301; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;2302; CHECK-SM80-NEXT: ret;2303;2304; CHECK-SM100-LABEL: reduce_and_i32(2305; CHECK-SM100: {2306; CHECK-SM100-NEXT: .reg .b32 %r<16>;2307; CHECK-SM100-NEXT: .reg .b64 %rd<5>;2308; CHECK-SM100-EMPTY:2309; CHECK-SM100-NEXT: // %bb.0:2310; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_and_i32_param_0+16];2311; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_and_i32_param_0];2312; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;2313; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;2314; CHECK-SM100-NEXT: and.b32 %r5, %r4, %r2;2315; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;2316; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;2317; CHECK-SM100-NEXT: and.b32 %r10, %r9, %r7;2318; CHECK-SM100-NEXT: and.b32 %r11, %r10, %r5;2319; CHECK-SM100-NEXT: and.b32 %r12, %r3, %r1;2320; CHECK-SM100-NEXT: and.b32 %r13, %r8, %r6;2321; CHECK-SM100-NEXT: and.b32 %r14, %r13, %r12;2322; CHECK-SM100-NEXT: and.b32 %r15, %r14, %r11;2323; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;2324; CHECK-SM100-NEXT: ret;2325 %res = call i32 @llvm.vector.reduce.and(<8 x i32> %in)2326 ret i32 %res2327}2328 2329define i32 @reduce_and_i32_nonpow2(<7 x i32> %in) {2330; CHECK-LABEL: reduce_and_i32_nonpow2(2331; CHECK: {2332; CHECK-NEXT: .reg .b32 %r<14>;2333; CHECK-EMPTY:2334; CHECK-NEXT: // %bb.0:2335; CHECK-NEXT: ld.param.b32 %r7, [reduce_and_i32_nonpow2_param_0+24];2336; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_and_i32_nonpow2_param_0+16];2337; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_and_i32_nonpow2_param_0];2338; CHECK-NEXT: and.b32 %r8, %r3, %r7;2339; CHECK-NEXT: and.b32 %r9, %r1, %r5;2340; CHECK-NEXT: and.b32 %r10, %r9, %r8;2341; CHECK-NEXT: and.b32 %r11, %r2, %r6;2342; CHECK-NEXT: and.b32 %r12, %r11, %r4;2343; CHECK-NEXT: and.b32 %r13, %r10, %r12;2344; CHECK-NEXT: st.param.b32 [func_retval0], %r13;2345; CHECK-NEXT: ret;2346 %res = call i32 @llvm.vector.reduce.and(<7 x i32> %in)2347 ret i32 %res2348}2349 2350define i16 @reduce_or_i16(<8 x i16> %in) {2351; CHECK-LABEL: reduce_or_i16(2352; CHECK: {2353; CHECK-NEXT: .reg .b16 %rs<4>;2354; CHECK-NEXT: .reg .b32 %r<9>;2355; CHECK-EMPTY:2356; CHECK-NEXT: // %bb.0:2357; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_or_i16_param_0];2358; CHECK-NEXT: or.b32 %r5, %r2, %r4;2359; CHECK-NEXT: or.b32 %r6, %r1, %r3;2360; CHECK-NEXT: or.b32 %r7, %r6, %r5;2361; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;2362; CHECK-NEXT: or.b16 %rs3, %rs1, %rs2;2363; CHECK-NEXT: cvt.u32.u16 %r8, %rs3;2364; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2365; CHECK-NEXT: ret;2366 %res = call i16 @llvm.vector.reduce.or(<8 x i16> %in)2367 ret i16 %res2368}2369 2370define i16 @reduce_or_i16_nonpow2(<7 x i16> %in) {2371; CHECK-LABEL: reduce_or_i16_nonpow2(2372; CHECK: {2373; CHECK-NEXT: .reg .b16 %rs<12>;2374; CHECK-NEXT: .reg .b32 %r<9>;2375; CHECK-EMPTY:2376; CHECK-NEXT: // %bb.0:2377; CHECK-NEXT: ld.param.b32 %r1, [reduce_or_i16_nonpow2_param_0+8];2378; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;2379; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_or_i16_nonpow2_param_0];2380; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;2381; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;2382; CHECK-NEXT: ld.param.b16 %rs7, [reduce_or_i16_nonpow2_param_0+12];2383; CHECK-NEXT: or.b32 %r4, %r2, %r1;2384; CHECK-NEXT: mov.b16 %rs8, 0;2385; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};2386; CHECK-NEXT: or.b32 %r6, %r3, %r5;2387; CHECK-NEXT: or.b32 %r7, %r4, %r6;2388; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;2389; CHECK-NEXT: or.b16 %rs11, %rs9, %rs10;2390; CHECK-NEXT: cvt.u32.u16 %r8, %rs11;2391; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2392; CHECK-NEXT: ret;2393 %res = call i16 @llvm.vector.reduce.or(<7 x i16> %in)2394 ret i16 %res2395}2396 2397define i32 @reduce_or_i32(<8 x i32> %in) {2398; CHECK-SM80-LABEL: reduce_or_i32(2399; CHECK-SM80: {2400; CHECK-SM80-NEXT: .reg .b32 %r<16>;2401; CHECK-SM80-EMPTY:2402; CHECK-SM80-NEXT: // %bb.0:2403; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_or_i32_param_0+16];2404; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_or_i32_param_0];2405; CHECK-SM80-NEXT: or.b32 %r9, %r4, %r8;2406; CHECK-SM80-NEXT: or.b32 %r10, %r2, %r6;2407; CHECK-SM80-NEXT: or.b32 %r11, %r10, %r9;2408; CHECK-SM80-NEXT: or.b32 %r12, %r3, %r7;2409; CHECK-SM80-NEXT: or.b32 %r13, %r1, %r5;2410; CHECK-SM80-NEXT: or.b32 %r14, %r13, %r12;2411; CHECK-SM80-NEXT: or.b32 %r15, %r14, %r11;2412; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;2413; CHECK-SM80-NEXT: ret;2414;2415; CHECK-SM100-LABEL: reduce_or_i32(2416; CHECK-SM100: {2417; CHECK-SM100-NEXT: .reg .b32 %r<16>;2418; CHECK-SM100-NEXT: .reg .b64 %rd<5>;2419; CHECK-SM100-EMPTY:2420; CHECK-SM100-NEXT: // %bb.0:2421; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_or_i32_param_0+16];2422; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_or_i32_param_0];2423; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;2424; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;2425; CHECK-SM100-NEXT: or.b32 %r5, %r4, %r2;2426; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;2427; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;2428; CHECK-SM100-NEXT: or.b32 %r10, %r9, %r7;2429; CHECK-SM100-NEXT: or.b32 %r11, %r10, %r5;2430; CHECK-SM100-NEXT: or.b32 %r12, %r3, %r1;2431; CHECK-SM100-NEXT: or.b32 %r13, %r8, %r6;2432; CHECK-SM100-NEXT: or.b32 %r14, %r13, %r12;2433; CHECK-SM100-NEXT: or.b32 %r15, %r14, %r11;2434; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;2435; CHECK-SM100-NEXT: ret;2436 %res = call i32 @llvm.vector.reduce.or(<8 x i32> %in)2437 ret i32 %res2438}2439 2440define i32 @reduce_or_i32_nonpow2(<7 x i32> %in) {2441; CHECK-LABEL: reduce_or_i32_nonpow2(2442; CHECK: {2443; CHECK-NEXT: .reg .b32 %r<14>;2444; CHECK-EMPTY:2445; CHECK-NEXT: // %bb.0:2446; CHECK-NEXT: ld.param.b32 %r7, [reduce_or_i32_nonpow2_param_0+24];2447; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_or_i32_nonpow2_param_0+16];2448; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_or_i32_nonpow2_param_0];2449; CHECK-NEXT: or.b32 %r8, %r3, %r7;2450; CHECK-NEXT: or.b32 %r9, %r1, %r5;2451; CHECK-NEXT: or.b32 %r10, %r9, %r8;2452; CHECK-NEXT: or.b32 %r11, %r2, %r6;2453; CHECK-NEXT: or.b32 %r12, %r11, %r4;2454; CHECK-NEXT: or.b32 %r13, %r10, %r12;2455; CHECK-NEXT: st.param.b32 [func_retval0], %r13;2456; CHECK-NEXT: ret;2457 %res = call i32 @llvm.vector.reduce.or(<7 x i32> %in)2458 ret i32 %res2459}2460 2461define i16 @reduce_xor_i16(<8 x i16> %in) {2462; CHECK-LABEL: reduce_xor_i16(2463; CHECK: {2464; CHECK-NEXT: .reg .b16 %rs<4>;2465; CHECK-NEXT: .reg .b32 %r<9>;2466; CHECK-EMPTY:2467; CHECK-NEXT: // %bb.0:2468; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_xor_i16_param_0];2469; CHECK-NEXT: xor.b32 %r5, %r2, %r4;2470; CHECK-NEXT: xor.b32 %r6, %r1, %r3;2471; CHECK-NEXT: xor.b32 %r7, %r6, %r5;2472; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r7;2473; CHECK-NEXT: xor.b16 %rs3, %rs1, %rs2;2474; CHECK-NEXT: cvt.u32.u16 %r8, %rs3;2475; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2476; CHECK-NEXT: ret;2477 %res = call i16 @llvm.vector.reduce.xor(<8 x i16> %in)2478 ret i16 %res2479}2480 2481define i16 @reduce_xor_i16_nonpow2(<7 x i16> %in) {2482; CHECK-LABEL: reduce_xor_i16_nonpow2(2483; CHECK: {2484; CHECK-NEXT: .reg .b16 %rs<12>;2485; CHECK-NEXT: .reg .b32 %r<9>;2486; CHECK-EMPTY:2487; CHECK-NEXT: // %bb.0:2488; CHECK-NEXT: ld.param.b32 %r1, [reduce_xor_i16_nonpow2_param_0+8];2489; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;2490; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [reduce_xor_i16_nonpow2_param_0];2491; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r3;2492; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;2493; CHECK-NEXT: ld.param.b16 %rs7, [reduce_xor_i16_nonpow2_param_0+12];2494; CHECK-NEXT: xor.b32 %r4, %r2, %r1;2495; CHECK-NEXT: mov.b16 %rs8, 0;2496; CHECK-NEXT: mov.b32 %r5, {%rs7, %rs8};2497; CHECK-NEXT: xor.b32 %r6, %r3, %r5;2498; CHECK-NEXT: xor.b32 %r7, %r4, %r6;2499; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r7;2500; CHECK-NEXT: xor.b16 %rs11, %rs9, %rs10;2501; CHECK-NEXT: cvt.u32.u16 %r8, %rs11;2502; CHECK-NEXT: st.param.b32 [func_retval0], %r8;2503; CHECK-NEXT: ret;2504 %res = call i16 @llvm.vector.reduce.xor(<7 x i16> %in)2505 ret i16 %res2506}2507 2508define i32 @reduce_xor_i32(<8 x i32> %in) {2509; CHECK-SM80-LABEL: reduce_xor_i32(2510; CHECK-SM80: {2511; CHECK-SM80-NEXT: .reg .b32 %r<16>;2512; CHECK-SM80-EMPTY:2513; CHECK-SM80-NEXT: // %bb.0:2514; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_xor_i32_param_0+16];2515; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_xor_i32_param_0];2516; CHECK-SM80-NEXT: xor.b32 %r9, %r4, %r8;2517; CHECK-SM80-NEXT: xor.b32 %r10, %r2, %r6;2518; CHECK-SM80-NEXT: xor.b32 %r11, %r10, %r9;2519; CHECK-SM80-NEXT: xor.b32 %r12, %r3, %r7;2520; CHECK-SM80-NEXT: xor.b32 %r13, %r1, %r5;2521; CHECK-SM80-NEXT: xor.b32 %r14, %r13, %r12;2522; CHECK-SM80-NEXT: xor.b32 %r15, %r14, %r11;2523; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;2524; CHECK-SM80-NEXT: ret;2525;2526; CHECK-SM100-LABEL: reduce_xor_i32(2527; CHECK-SM100: {2528; CHECK-SM100-NEXT: .reg .b32 %r<16>;2529; CHECK-SM100-NEXT: .reg .b64 %rd<5>;2530; CHECK-SM100-EMPTY:2531; CHECK-SM100-NEXT: // %bb.0:2532; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_xor_i32_param_0+16];2533; CHECK-SM100-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_xor_i32_param_0];2534; CHECK-SM100-NEXT: mov.b64 {%r1, %r2}, %rd4;2535; CHECK-SM100-NEXT: mov.b64 {%r3, %r4}, %rd2;2536; CHECK-SM100-NEXT: xor.b32 %r5, %r4, %r2;2537; CHECK-SM100-NEXT: mov.b64 {%r6, %r7}, %rd3;2538; CHECK-SM100-NEXT: mov.b64 {%r8, %r9}, %rd1;2539; CHECK-SM100-NEXT: xor.b32 %r10, %r9, %r7;2540; CHECK-SM100-NEXT: xor.b32 %r11, %r10, %r5;2541; CHECK-SM100-NEXT: xor.b32 %r12, %r3, %r1;2542; CHECK-SM100-NEXT: xor.b32 %r13, %r8, %r6;2543; CHECK-SM100-NEXT: xor.b32 %r14, %r13, %r12;2544; CHECK-SM100-NEXT: xor.b32 %r15, %r14, %r11;2545; CHECK-SM100-NEXT: st.param.b32 [func_retval0], %r15;2546; CHECK-SM100-NEXT: ret;2547 %res = call i32 @llvm.vector.reduce.xor(<8 x i32> %in)2548 ret i32 %res2549}2550 2551define i32 @reduce_xor_i32_nonpow2(<7 x i32> %in) {2552; CHECK-LABEL: reduce_xor_i32_nonpow2(2553; CHECK: {2554; CHECK-NEXT: .reg .b32 %r<14>;2555; CHECK-EMPTY:2556; CHECK-NEXT: // %bb.0:2557; CHECK-NEXT: ld.param.b32 %r7, [reduce_xor_i32_nonpow2_param_0+24];2558; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [reduce_xor_i32_nonpow2_param_0+16];2559; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_xor_i32_nonpow2_param_0];2560; CHECK-NEXT: xor.b32 %r8, %r3, %r7;2561; CHECK-NEXT: xor.b32 %r9, %r1, %r5;2562; CHECK-NEXT: xor.b32 %r10, %r9, %r8;2563; CHECK-NEXT: xor.b32 %r11, %r2, %r6;2564; CHECK-NEXT: xor.b32 %r12, %r11, %r4;2565; CHECK-NEXT: xor.b32 %r13, %r10, %r12;2566; CHECK-NEXT: st.param.b32 [func_retval0], %r13;2567; CHECK-NEXT: ret;2568 %res = call i32 @llvm.vector.reduce.xor(<7 x i32> %in)2569 ret i32 %res2570}2571