319 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 | %ptxas-verify %}4 5declare {i32, i1} @llvm.nvvm.shfl.down.i32p(i32, i32, i32)6declare {float, i1} @llvm.nvvm.shfl.down.f32p(float, i32, i32)7declare {i32, i1} @llvm.nvvm.shfl.up.i32p(i32, i32, i32)8declare {float, i1} @llvm.nvvm.shfl.up.f32p(float, i32, i32)9declare {i32, i1} @llvm.nvvm.shfl.bfly.i32p(i32, i32, i32)10declare {float, i1} @llvm.nvvm.shfl.bfly.f32p(float, i32, i32)11declare {i32, i1} @llvm.nvvm.shfl.idx.i32p(i32, i32, i32)12declare {float, i1} @llvm.nvvm.shfl.idx.f32p(float, i32, i32)13 14define {i32, i1} @shfl_i32_rrr(i32 %a, i32 %b, i32 %c) {15; CHECK-LABEL: shfl_i32_rrr(16; CHECK: {17; CHECK-NEXT: .reg .pred %p<2>;18; CHECK-NEXT: .reg .b16 %rs<2>;19; CHECK-NEXT: .reg .b32 %r<5>;20; CHECK-EMPTY:21; CHECK-NEXT: // %bb.0:22; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_rrr_param_0];23; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_rrr_param_1];24; CHECK-NEXT: ld.param.b32 %r3, [shfl_i32_rrr_param_2];25; CHECK-NEXT: shfl.down.b32 %r4|%p1, %r1, %r2, %r3;26; CHECK-NEXT: st.param.b32 [func_retval0], %r4;27; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;28; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;29; CHECK-NEXT: ret;30 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)31 ret {i32, i1} %val32}33 34define {i32, i1} @shfl_i32_irr(i32 %a, i32 %b, i32 %c) {35; CHECK-LABEL: shfl_i32_irr(36; CHECK: {37; CHECK-NEXT: .reg .pred %p<2>;38; CHECK-NEXT: .reg .b16 %rs<2>;39; CHECK-NEXT: .reg .b32 %r<5>;40; CHECK-EMPTY:41; CHECK-NEXT: // %bb.0:42; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_irr_param_0];43; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_irr_param_1];44; CHECK-NEXT: ld.param.b32 %r3, [shfl_i32_irr_param_2];45; CHECK-NEXT: shfl.down.b32 %r4|%p1, %r1, %r2, %r3;46; CHECK-NEXT: st.param.b32 [func_retval0], %r4;47; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;48; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;49; CHECK-NEXT: ret;50 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)51 ret {i32, i1} %val52}53 54define {i32, i1} @shfl_i32_rri(i32 %a, i32 %b) {55; CHECK-LABEL: shfl_i32_rri(56; CHECK: {57; CHECK-NEXT: .reg .pred %p<2>;58; CHECK-NEXT: .reg .b16 %rs<2>;59; CHECK-NEXT: .reg .b32 %r<4>;60; CHECK-EMPTY:61; CHECK-NEXT: // %bb.0:62; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_rri_param_0];63; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_rri_param_1];64; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, %r2, 1;65; CHECK-NEXT: st.param.b32 [func_retval0], %r3;66; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;67; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;68; CHECK-NEXT: ret;69 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 1)70 ret {i32, i1} %val71}72 73define {i32, i1} @shfl_i32_iri(i32 %a, i32 %b) {74; CHECK-LABEL: shfl_i32_iri(75; CHECK: {76; CHECK-NEXT: .reg .pred %p<2>;77; CHECK-NEXT: .reg .b16 %rs<2>;78; CHECK-NEXT: .reg .b32 %r<4>;79; CHECK-EMPTY:80; CHECK-NEXT: // %bb.0:81; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_iri_param_0];82; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_iri_param_1];83; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, %r2, 2;84; CHECK-NEXT: st.param.b32 [func_retval0], %r3;85; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;86; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;87; CHECK-NEXT: ret;88 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 2)89 ret {i32, i1} %val90}91 92define {i32, i1} @shfl_i32_rir(i32 %a, i32 %c) {93; CHECK-LABEL: shfl_i32_rir(94; CHECK: {95; CHECK-NEXT: .reg .pred %p<2>;96; CHECK-NEXT: .reg .b16 %rs<2>;97; CHECK-NEXT: .reg .b32 %r<4>;98; CHECK-EMPTY:99; CHECK-NEXT: // %bb.0:100; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_rir_param_0];101; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_rir_param_1];102; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, 1, %r2;103; CHECK-NEXT: st.param.b32 [func_retval0], %r3;104; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;105; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;106; CHECK-NEXT: ret;107 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 1, i32 %c)108 ret {i32, i1} %val109}110 111define {i32, i1} @shfl_i32_iir(i32 %a, i32 %c) {112; CHECK-LABEL: shfl_i32_iir(113; CHECK: {114; CHECK-NEXT: .reg .pred %p<2>;115; CHECK-NEXT: .reg .b16 %rs<2>;116; CHECK-NEXT: .reg .b32 %r<4>;117; CHECK-EMPTY:118; CHECK-NEXT: // %bb.0:119; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_iir_param_0];120; CHECK-NEXT: ld.param.b32 %r2, [shfl_i32_iir_param_1];121; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, 2, %r2;122; CHECK-NEXT: st.param.b32 [func_retval0], %r3;123; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;124; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;125; CHECK-NEXT: ret;126 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 2, i32 %c)127 ret {i32, i1} %val128}129 130define {i32, i1} @shfl_i32_rii(i32 %a) {131; CHECK-LABEL: shfl_i32_rii(132; CHECK: {133; CHECK-NEXT: .reg .pred %p<2>;134; CHECK-NEXT: .reg .b16 %rs<2>;135; CHECK-NEXT: .reg .b32 %r<3>;136; CHECK-EMPTY:137; CHECK-NEXT: // %bb.0:138; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_rii_param_0];139; CHECK-NEXT: shfl.down.b32 %r2|%p1, %r1, 1, 2;140; CHECK-NEXT: st.param.b32 [func_retval0], %r2;141; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;142; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;143; CHECK-NEXT: ret;144 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 1, i32 2)145 ret {i32, i1} %val146}147 148define {i32, i1} @shfl_i32_iii(i32 %a, i32 %b) {149; CHECK-LABEL: shfl_i32_iii(150; CHECK: {151; CHECK-NEXT: .reg .pred %p<2>;152; CHECK-NEXT: .reg .b16 %rs<2>;153; CHECK-NEXT: .reg .b32 %r<3>;154; CHECK-EMPTY:155; CHECK-NEXT: // %bb.0:156; CHECK-NEXT: ld.param.b32 %r1, [shfl_i32_iii_param_0];157; CHECK-NEXT: shfl.down.b32 %r2|%p1, %r1, 2, 3;158; CHECK-NEXT: st.param.b32 [func_retval0], %r2;159; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;160; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;161; CHECK-NEXT: ret;162 %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 2, i32 3)163 ret {i32, i1} %val164}165 166;; Same intrinsics, but for float167 168define {float, i1} @shfl_f32_rrr(float %a, i32 %b, i32 %c) {169; CHECK-LABEL: shfl_f32_rrr(170; CHECK: {171; CHECK-NEXT: .reg .pred %p<2>;172; CHECK-NEXT: .reg .b16 %rs<2>;173; CHECK-NEXT: .reg .b32 %r<5>;174; CHECK-EMPTY:175; CHECK-NEXT: // %bb.0:176; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_rrr_param_0];177; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_rrr_param_1];178; CHECK-NEXT: ld.param.b32 %r3, [shfl_f32_rrr_param_2];179; CHECK-NEXT: shfl.down.b32 %r4|%p1, %r1, %r2, %r3;180; CHECK-NEXT: st.param.b32 [func_retval0], %r4;181; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;182; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;183; CHECK-NEXT: ret;184 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 %c)185 ret {float, i1} %val186}187 188define {float, i1} @shfl_f32_irr(float %a, i32 %b, i32 %c) {189; CHECK-LABEL: shfl_f32_irr(190; CHECK: {191; CHECK-NEXT: .reg .pred %p<2>;192; CHECK-NEXT: .reg .b16 %rs<2>;193; CHECK-NEXT: .reg .b32 %r<5>;194; CHECK-EMPTY:195; CHECK-NEXT: // %bb.0:196; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_irr_param_0];197; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_irr_param_1];198; CHECK-NEXT: ld.param.b32 %r3, [shfl_f32_irr_param_2];199; CHECK-NEXT: shfl.down.b32 %r4|%p1, %r1, %r2, %r3;200; CHECK-NEXT: st.param.b32 [func_retval0], %r4;201; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;202; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;203; CHECK-NEXT: ret;204 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 %c)205 ret {float, i1} %val206}207 208define {float, i1} @shfl_f32_rri(float %a, i32 %b) {209; CHECK-LABEL: shfl_f32_rri(210; CHECK: {211; CHECK-NEXT: .reg .pred %p<2>;212; CHECK-NEXT: .reg .b16 %rs<2>;213; CHECK-NEXT: .reg .b32 %r<4>;214; CHECK-EMPTY:215; CHECK-NEXT: // %bb.0:216; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_rri_param_0];217; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_rri_param_1];218; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, %r2, 1;219; CHECK-NEXT: st.param.b32 [func_retval0], %r3;220; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;221; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;222; CHECK-NEXT: ret;223 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 1)224 ret {float, i1} %val225}226 227define {float, i1} @shfl_f32_iri(float %a, i32 %b) {228; CHECK-LABEL: shfl_f32_iri(229; CHECK: {230; CHECK-NEXT: .reg .pred %p<2>;231; CHECK-NEXT: .reg .b16 %rs<2>;232; CHECK-NEXT: .reg .b32 %r<4>;233; CHECK-EMPTY:234; CHECK-NEXT: // %bb.0:235; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_iri_param_0];236; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_iri_param_1];237; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, %r2, 2;238; CHECK-NEXT: st.param.b32 [func_retval0], %r3;239; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;240; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;241; CHECK-NEXT: ret;242 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 2)243 ret {float, i1} %val244}245 246define {float, i1} @shfl_f32_rir(float %a, i32 %c) {247; CHECK-LABEL: shfl_f32_rir(248; CHECK: {249; CHECK-NEXT: .reg .pred %p<2>;250; CHECK-NEXT: .reg .b16 %rs<2>;251; CHECK-NEXT: .reg .b32 %r<4>;252; CHECK-EMPTY:253; CHECK-NEXT: // %bb.0:254; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_rir_param_0];255; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_rir_param_1];256; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, 1, %r2;257; CHECK-NEXT: st.param.b32 [func_retval0], %r3;258; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;259; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;260; CHECK-NEXT: ret;261 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 1, i32 %c)262 ret {float, i1} %val263}264 265define {float, i1} @shfl_f32_iir(float %a, i32 %c) {266; CHECK-LABEL: shfl_f32_iir(267; CHECK: {268; CHECK-NEXT: .reg .pred %p<2>;269; CHECK-NEXT: .reg .b16 %rs<2>;270; CHECK-NEXT: .reg .b32 %r<4>;271; CHECK-EMPTY:272; CHECK-NEXT: // %bb.0:273; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_iir_param_0];274; CHECK-NEXT: ld.param.b32 %r2, [shfl_f32_iir_param_1];275; CHECK-NEXT: shfl.down.b32 %r3|%p1, %r1, 2, %r2;276; CHECK-NEXT: st.param.b32 [func_retval0], %r3;277; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;278; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;279; CHECK-NEXT: ret;280 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 2, i32 %c)281 ret {float, i1} %val282}283 284define {float, i1} @shfl_f32_rii(float %a) {285; CHECK-LABEL: shfl_f32_rii(286; CHECK: {287; CHECK-NEXT: .reg .pred %p<2>;288; CHECK-NEXT: .reg .b16 %rs<2>;289; CHECK-NEXT: .reg .b32 %r<3>;290; CHECK-EMPTY:291; CHECK-NEXT: // %bb.0:292; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_rii_param_0];293; CHECK-NEXT: shfl.down.b32 %r2|%p1, %r1, 1, 2;294; CHECK-NEXT: st.param.b32 [func_retval0], %r2;295; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;296; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;297; CHECK-NEXT: ret;298 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 1, i32 2)299 ret {float, i1} %val300}301 302define {float, i1} @shfl_f32_iii(float %a, i32 %b) {303; CHECK-LABEL: shfl_f32_iii(304; CHECK: {305; CHECK-NEXT: .reg .pred %p<2>;306; CHECK-NEXT: .reg .b16 %rs<2>;307; CHECK-NEXT: .reg .b32 %r<3>;308; CHECK-EMPTY:309; CHECK-NEXT: // %bb.0:310; CHECK-NEXT: ld.param.b32 %r1, [shfl_f32_iii_param_0];311; CHECK-NEXT: shfl.down.b32 %r2|%p1, %r1, 2, 3;312; CHECK-NEXT: st.param.b32 [func_retval0], %r2;313; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p1;314; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs1;315; CHECK-NEXT: ret;316 %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 2, i32 3)317 ret {float, i1} %val318}319