3236 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5; TODO: add i1, and <6 x i8> vector tests.6 7; TODO: add test for vectors that exceed 128-bit length8; Per https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#vectors9; vectors cannot exceed 128-bit in length, i.e., .v4.u64 is not allowed.10 11; TODO: generate PTX that preserves Concurrent Forward Progress12; for atomic operations to local statespace13; by generating atomic or volatile operations.14 15; TODO: design exposure for atomic operations on vector types.16 17; TODO: add weak,atomic,volatile,atomic volatile tests18; for .const and .param statespaces.19 20;; generic statespace21 22; generic23 24; TODO: make the lowering of this weak vector ops consistent with25; the ones of the next tests. This test lowers to a weak PTX26; vector op, but next test lowers to a vector PTX op.27define void @generic_2xi8(ptr %a) {28; CHECK-LABEL: generic_2xi8(29; CHECK: {30; CHECK-NEXT: .reg .b16 %rs<5>;31; CHECK-NEXT: .reg .b64 %rd<2>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xi8_param_0];35; CHECK-NEXT: ld.v2.b8 {%rs1, %rs2}, [%rd1];36; CHECK-NEXT: add.s16 %rs3, %rs2, 1;37; CHECK-NEXT: add.s16 %rs4, %rs1, 1;38; CHECK-NEXT: st.v2.b8 [%rd1], {%rs4, %rs3};39; CHECK-NEXT: ret;40 %a.load = load <2 x i8>, ptr %a41 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>42 store <2 x i8> %a.add, ptr %a43 ret void44}45 46; TODO: make the lowering of this weak vector ops consistent with47; the ones of the previous test. This test lowers to a weak48; PTX scalar op, but prior test lowers to a vector PTX op.49define void @generic_4xi8(ptr %a) {50; CHECK-LABEL: generic_4xi8(51; CHECK: {52; CHECK-NEXT: .reg .b16 %rs<9>;53; CHECK-NEXT: .reg .b32 %r<13>;54; CHECK-NEXT: .reg .b64 %rd<2>;55; CHECK-EMPTY:56; CHECK-NEXT: // %bb.0:57; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xi8_param_0];58; CHECK-NEXT: ld.b32 %r1, [%rd1];59; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;60; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;61; CHECK-NEXT: add.s16 %rs2, %rs1, 1;62; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;63; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;64; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;65; CHECK-NEXT: add.s16 %rs4, %rs3, 1;66; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;67; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;68; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;69; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;70; CHECK-NEXT: add.s16 %rs6, %rs5, 1;71; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;72; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;73; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;74; CHECK-NEXT: add.s16 %rs8, %rs7, 1;75; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;76; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;77; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;78; CHECK-NEXT: st.b32 [%rd1], %r12;79; CHECK-NEXT: ret;80 %a.load = load <4 x i8>, ptr %a81 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>82 store <4 x i8> %a.add, ptr %a83 ret void84}85 86define void @generic_8xi8(ptr %a) {87; CHECK-LABEL: generic_8xi8(88; CHECK: {89; CHECK-NEXT: .reg .b16 %rs<17>;90; CHECK-NEXT: .reg .b32 %r<25>;91; CHECK-NEXT: .reg .b64 %rd<2>;92; CHECK-EMPTY:93; CHECK-NEXT: // %bb.0:94; CHECK-NEXT: ld.param.b64 %rd1, [generic_8xi8_param_0];95; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];96; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;97; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;98; CHECK-NEXT: add.s16 %rs2, %rs1, 1;99; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;100; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;101; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;102; CHECK-NEXT: add.s16 %rs4, %rs3, 1;103; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;104; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;105; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;106; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;107; CHECK-NEXT: add.s16 %rs6, %rs5, 1;108; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;109; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;110; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;111; CHECK-NEXT: add.s16 %rs8, %rs7, 1;112; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;113; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;114; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;115; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;116; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;117; CHECK-NEXT: add.s16 %rs10, %rs9, 1;118; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;119; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;120; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;121; CHECK-NEXT: add.s16 %rs12, %rs11, 1;122; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;123; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;124; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;125; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;126; CHECK-NEXT: add.s16 %rs14, %rs13, 1;127; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;128; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;129; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;130; CHECK-NEXT: add.s16 %rs16, %rs15, 1;131; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;132; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;133; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;134; CHECK-NEXT: st.v2.b32 [%rd1], {%r24, %r13};135; CHECK-NEXT: ret;136 %a.load = load <8 x i8>, ptr %a137 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>138 store <8 x i8> %a.add, ptr %a139 ret void140}141 142define void @generic_16xi8(ptr %a) {143; CHECK-LABEL: generic_16xi8(144; CHECK: {145; CHECK-NEXT: .reg .b16 %rs<33>;146; CHECK-NEXT: .reg .b32 %r<49>;147; CHECK-NEXT: .reg .b64 %rd<2>;148; CHECK-EMPTY:149; CHECK-NEXT: // %bb.0:150; CHECK-NEXT: ld.param.b64 %rd1, [generic_16xi8_param_0];151; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];152; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;153; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;154; CHECK-NEXT: add.s16 %rs2, %rs1, 1;155; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;156; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;157; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;158; CHECK-NEXT: add.s16 %rs4, %rs3, 1;159; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;160; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;161; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;162; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;163; CHECK-NEXT: add.s16 %rs6, %rs5, 1;164; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;165; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;166; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;167; CHECK-NEXT: add.s16 %rs8, %rs7, 1;168; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;169; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;170; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;171; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;172; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;173; CHECK-NEXT: add.s16 %rs10, %rs9, 1;174; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;175; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;176; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;177; CHECK-NEXT: add.s16 %rs12, %rs11, 1;178; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;179; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;180; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;181; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;182; CHECK-NEXT: add.s16 %rs14, %rs13, 1;183; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;184; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;185; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;186; CHECK-NEXT: add.s16 %rs16, %rs15, 1;187; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;188; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;189; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;190; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;191; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;192; CHECK-NEXT: add.s16 %rs18, %rs17, 1;193; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;194; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;195; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;196; CHECK-NEXT: add.s16 %rs20, %rs19, 1;197; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;198; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;199; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;200; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;201; CHECK-NEXT: add.s16 %rs22, %rs21, 1;202; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;203; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;204; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;205; CHECK-NEXT: add.s16 %rs24, %rs23, 1;206; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;207; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;208; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;209; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;210; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;211; CHECK-NEXT: add.s16 %rs26, %rs25, 1;212; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;213; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;214; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;215; CHECK-NEXT: add.s16 %rs28, %rs27, 1;216; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;217; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;218; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;219; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;220; CHECK-NEXT: add.s16 %rs30, %rs29, 1;221; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;222; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;223; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;224; CHECK-NEXT: add.s16 %rs32, %rs31, 1;225; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;226; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;227; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;228; CHECK-NEXT: st.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};229; CHECK-NEXT: ret;230 %a.load = load <16 x i8>, ptr %a231 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>232 store <16 x i8> %a.add, ptr %a233 ret void234}235 236define void @generic_2xi16(ptr %a) {237; CHECK-LABEL: generic_2xi16(238; CHECK: {239; CHECK-NEXT: .reg .b16 %rs<5>;240; CHECK-NEXT: .reg .b64 %rd<2>;241; CHECK-EMPTY:242; CHECK-NEXT: // %bb.0:243; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xi16_param_0];244; CHECK-NEXT: ld.v2.b16 {%rs1, %rs2}, [%rd1];245; CHECK-NEXT: add.s16 %rs3, %rs2, 1;246; CHECK-NEXT: add.s16 %rs4, %rs1, 1;247; CHECK-NEXT: st.v2.b16 [%rd1], {%rs4, %rs3};248; CHECK-NEXT: ret;249 %a.load = load <2 x i16>, ptr %a250 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>251 store <2 x i16> %a.add, ptr %a252 ret void253}254 255define void @generic_4xi16(ptr %a) {256; CHECK-LABEL: generic_4xi16(257; CHECK: {258; CHECK-NEXT: .reg .b16 %rs<9>;259; CHECK-NEXT: .reg .b64 %rd<2>;260; CHECK-EMPTY:261; CHECK-NEXT: // %bb.0:262; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xi16_param_0];263; CHECK-NEXT: ld.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];264; CHECK-NEXT: add.s16 %rs5, %rs4, 1;265; CHECK-NEXT: add.s16 %rs6, %rs3, 1;266; CHECK-NEXT: add.s16 %rs7, %rs2, 1;267; CHECK-NEXT: add.s16 %rs8, %rs1, 1;268; CHECK-NEXT: st.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};269; CHECK-NEXT: ret;270 %a.load = load <4 x i16>, ptr %a271 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>272 store <4 x i16> %a.add, ptr %a273 ret void274}275 276define void @generic_8xi16(ptr %a) {277; CHECK-LABEL: generic_8xi16(278; CHECK: {279; CHECK-NEXT: .reg .b16 %rs<17>;280; CHECK-NEXT: .reg .b32 %r<9>;281; CHECK-NEXT: .reg .b64 %rd<2>;282; CHECK-EMPTY:283; CHECK-NEXT: // %bb.0:284; CHECK-NEXT: ld.param.b64 %rd1, [generic_8xi16_param_0];285; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];286; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;287; CHECK-NEXT: add.s16 %rs3, %rs2, 1;288; CHECK-NEXT: add.s16 %rs4, %rs1, 1;289; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};290; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;291; CHECK-NEXT: add.s16 %rs7, %rs6, 1;292; CHECK-NEXT: add.s16 %rs8, %rs5, 1;293; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};294; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;295; CHECK-NEXT: add.s16 %rs11, %rs10, 1;296; CHECK-NEXT: add.s16 %rs12, %rs9, 1;297; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};298; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;299; CHECK-NEXT: add.s16 %rs15, %rs14, 1;300; CHECK-NEXT: add.s16 %rs16, %rs13, 1;301; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};302; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};303; CHECK-NEXT: ret;304 %a.load = load <8 x i16>, ptr %a305 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>306 store <8 x i16> %a.add, ptr %a307 ret void308}309 310define void @generic_2xi32(ptr %a) {311; CHECK-LABEL: generic_2xi32(312; CHECK: {313; CHECK-NEXT: .reg .b32 %r<5>;314; CHECK-NEXT: .reg .b64 %rd<2>;315; CHECK-EMPTY:316; CHECK-NEXT: // %bb.0:317; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xi32_param_0];318; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];319; CHECK-NEXT: add.s32 %r3, %r2, 1;320; CHECK-NEXT: add.s32 %r4, %r1, 1;321; CHECK-NEXT: st.v2.b32 [%rd1], {%r4, %r3};322; CHECK-NEXT: ret;323 %a.load = load <2 x i32>, ptr %a324 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>325 store <2 x i32> %a.add, ptr %a326 ret void327}328 329define void @generic_4xi32(ptr %a) {330; CHECK-LABEL: generic_4xi32(331; CHECK: {332; CHECK-NEXT: .reg .b32 %r<9>;333; CHECK-NEXT: .reg .b64 %rd<2>;334; CHECK-EMPTY:335; CHECK-NEXT: // %bb.0:336; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xi32_param_0];337; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];338; CHECK-NEXT: add.s32 %r5, %r4, 1;339; CHECK-NEXT: add.s32 %r6, %r3, 1;340; CHECK-NEXT: add.s32 %r7, %r2, 1;341; CHECK-NEXT: add.s32 %r8, %r1, 1;342; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};343; CHECK-NEXT: ret;344 %a.load = load <4 x i32>, ptr %a345 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>346 store <4 x i32> %a.add, ptr %a347 ret void348}349 350define void @generic_2xi64(ptr %a) {351; CHECK-LABEL: generic_2xi64(352; CHECK: {353; CHECK-NEXT: .reg .b64 %rd<6>;354; CHECK-EMPTY:355; CHECK-NEXT: // %bb.0:356; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xi64_param_0];357; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];358; CHECK-NEXT: add.s64 %rd4, %rd3, 1;359; CHECK-NEXT: add.s64 %rd5, %rd2, 1;360; CHECK-NEXT: st.v2.b64 [%rd1], {%rd5, %rd4};361; CHECK-NEXT: ret;362 %a.load = load <2 x i64>, ptr %a363 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>364 store <2 x i64> %a.add, ptr %a365 ret void366}367 368define void @generic_2xfloat(ptr %a) {369; CHECK-LABEL: generic_2xfloat(370; CHECK: {371; CHECK-NEXT: .reg .b32 %r<5>;372; CHECK-NEXT: .reg .b64 %rd<2>;373; CHECK-EMPTY:374; CHECK-NEXT: // %bb.0:375; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xfloat_param_0];376; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];377; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;378; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;379; CHECK-NEXT: st.v2.b32 [%rd1], {%r4, %r3};380; CHECK-NEXT: ret;381 %a.load = load <2 x float>, ptr %a382 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>383 store <2 x float> %a.add, ptr %a384 ret void385}386 387define void @generic_4xfloat(ptr %a) {388; CHECK-LABEL: generic_4xfloat(389; CHECK: {390; CHECK-NEXT: .reg .b32 %r<9>;391; CHECK-NEXT: .reg .b64 %rd<2>;392; CHECK-EMPTY:393; CHECK-NEXT: // %bb.0:394; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xfloat_param_0];395; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];396; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;397; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;398; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;399; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;400; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};401; CHECK-NEXT: ret;402 %a.load = load <4 x float>, ptr %a403 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>404 store <4 x float> %a.add, ptr %a405 ret void406}407 408define void @generic_2xdouble(ptr %a) {409; CHECK-LABEL: generic_2xdouble(410; CHECK: {411; CHECK-NEXT: .reg .b64 %rd<6>;412; CHECK-EMPTY:413; CHECK-NEXT: // %bb.0:414; CHECK-NEXT: ld.param.b64 %rd1, [generic_2xdouble_param_0];415; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];416; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;417; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;418; CHECK-NEXT: st.v2.b64 [%rd1], {%rd5, %rd4};419; CHECK-NEXT: ret;420 %a.load = load <2 x double>, ptr %a421 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>422 store <2 x double> %a.add, ptr %a423 ret void424}425 426; generic_volatile427 428; TODO: volatile, atomic, and volatile atomic memory operations on vector types.429; Currently, LLVM:430; - does not allow atomic operations on vectors.431; - it allows volatile operations but not clear what that means.432; Following both semantics make sense in general and PTX supports both:433; - volatile/atomic/volatile atomic applies to the whole vector434; - volatile/atomic/volatile atomic applies elementwise435; Actions required:436; - clarify LLVM semantics for volatile on vectors and align the NVPTX backend with those437; Below tests show that the current implementation picks the semantics in an inconsistent way438; * volatile <2 x i8> lowers to "elementwise volatile"439; * <4 x i8> lowers to "full vector volatile"440; - provide support for vector atomics, e.g., by extending LLVM IR or via intrinsics441; - update tests in load-store-sm70.ll as well.442 443; TODO: make this operation consistent with the one for <4 x i8>444; This operation lowers to a "element wise volatile PTX operation".445define void @generic_volatile_2xi8(ptr %a) {446; CHECK-LABEL: generic_volatile_2xi8(447; CHECK: {448; CHECK-NEXT: .reg .b16 %rs<5>;449; CHECK-NEXT: .reg .b64 %rd<2>;450; CHECK-EMPTY:451; CHECK-NEXT: // %bb.0:452; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xi8_param_0];453; CHECK-NEXT: ld.volatile.v2.b8 {%rs1, %rs2}, [%rd1];454; CHECK-NEXT: add.s16 %rs3, %rs2, 1;455; CHECK-NEXT: add.s16 %rs4, %rs1, 1;456; CHECK-NEXT: st.volatile.v2.b8 [%rd1], {%rs4, %rs3};457; CHECK-NEXT: ret;458 %a.load = load volatile <2 x i8>, ptr %a459 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>460 store volatile <2 x i8> %a.add, ptr %a461 ret void462}463 464; TODO: make this operation consistent with the one for <2 x i8>465; This operation lowers to a "full vector volatile PTX operation".466define void @generic_volatile_4xi8(ptr %a) {467; CHECK-LABEL: generic_volatile_4xi8(468; CHECK: {469; CHECK-NEXT: .reg .b16 %rs<9>;470; CHECK-NEXT: .reg .b32 %r<13>;471; CHECK-NEXT: .reg .b64 %rd<2>;472; CHECK-EMPTY:473; CHECK-NEXT: // %bb.0:474; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xi8_param_0];475; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];476; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;477; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;478; CHECK-NEXT: add.s16 %rs2, %rs1, 1;479; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;480; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;481; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;482; CHECK-NEXT: add.s16 %rs4, %rs3, 1;483; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;484; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;485; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;486; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;487; CHECK-NEXT: add.s16 %rs6, %rs5, 1;488; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;489; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;490; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;491; CHECK-NEXT: add.s16 %rs8, %rs7, 1;492; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;493; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;494; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;495; CHECK-NEXT: st.volatile.b32 [%rd1], %r12;496; CHECK-NEXT: ret;497 %a.load = load volatile <4 x i8>, ptr %a498 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>499 store volatile <4 x i8> %a.add, ptr %a500 ret void501}502 503define void @generic_volatile_8xi8(ptr %a) {504; CHECK-LABEL: generic_volatile_8xi8(505; CHECK: {506; CHECK-NEXT: .reg .b16 %rs<17>;507; CHECK-NEXT: .reg .b32 %r<25>;508; CHECK-NEXT: .reg .b64 %rd<2>;509; CHECK-EMPTY:510; CHECK-NEXT: // %bb.0:511; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_8xi8_param_0];512; CHECK-NEXT: ld.volatile.v2.b32 {%r1, %r2}, [%rd1];513; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;514; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;515; CHECK-NEXT: add.s16 %rs2, %rs1, 1;516; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;517; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;518; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;519; CHECK-NEXT: add.s16 %rs4, %rs3, 1;520; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;521; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;522; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;523; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;524; CHECK-NEXT: add.s16 %rs6, %rs5, 1;525; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;526; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;527; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;528; CHECK-NEXT: add.s16 %rs8, %rs7, 1;529; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;530; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;531; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;532; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;533; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;534; CHECK-NEXT: add.s16 %rs10, %rs9, 1;535; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;536; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;537; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;538; CHECK-NEXT: add.s16 %rs12, %rs11, 1;539; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;540; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;541; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;542; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;543; CHECK-NEXT: add.s16 %rs14, %rs13, 1;544; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;545; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;546; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;547; CHECK-NEXT: add.s16 %rs16, %rs15, 1;548; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;549; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;550; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;551; CHECK-NEXT: st.volatile.v2.b32 [%rd1], {%r24, %r13};552; CHECK-NEXT: ret;553 %a.load = load volatile <8 x i8>, ptr %a554 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>555 store volatile <8 x i8> %a.add, ptr %a556 ret void557}558 559define void @generic_volatile_16xi8(ptr %a) {560; CHECK-LABEL: generic_volatile_16xi8(561; CHECK: {562; CHECK-NEXT: .reg .b16 %rs<33>;563; CHECK-NEXT: .reg .b32 %r<49>;564; CHECK-NEXT: .reg .b64 %rd<2>;565; CHECK-EMPTY:566; CHECK-NEXT: // %bb.0:567; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_16xi8_param_0];568; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];569; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;570; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;571; CHECK-NEXT: add.s16 %rs2, %rs1, 1;572; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;573; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;574; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;575; CHECK-NEXT: add.s16 %rs4, %rs3, 1;576; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;577; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;578; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;579; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;580; CHECK-NEXT: add.s16 %rs6, %rs5, 1;581; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;582; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;583; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;584; CHECK-NEXT: add.s16 %rs8, %rs7, 1;585; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;586; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;587; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;588; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;589; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;590; CHECK-NEXT: add.s16 %rs10, %rs9, 1;591; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;592; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;593; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;594; CHECK-NEXT: add.s16 %rs12, %rs11, 1;595; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;596; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;597; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;598; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;599; CHECK-NEXT: add.s16 %rs14, %rs13, 1;600; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;601; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;602; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;603; CHECK-NEXT: add.s16 %rs16, %rs15, 1;604; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;605; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;606; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;607; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;608; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;609; CHECK-NEXT: add.s16 %rs18, %rs17, 1;610; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;611; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;612; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;613; CHECK-NEXT: add.s16 %rs20, %rs19, 1;614; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;615; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;616; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;617; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;618; CHECK-NEXT: add.s16 %rs22, %rs21, 1;619; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;620; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;621; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;622; CHECK-NEXT: add.s16 %rs24, %rs23, 1;623; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;624; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;625; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;626; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;627; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;628; CHECK-NEXT: add.s16 %rs26, %rs25, 1;629; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;630; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;631; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;632; CHECK-NEXT: add.s16 %rs28, %rs27, 1;633; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;634; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;635; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;636; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;637; CHECK-NEXT: add.s16 %rs30, %rs29, 1;638; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;639; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;640; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;641; CHECK-NEXT: add.s16 %rs32, %rs31, 1;642; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;643; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;644; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;645; CHECK-NEXT: st.volatile.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};646; CHECK-NEXT: ret;647 %a.load = load volatile <16 x i8>, ptr %a648 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>649 store volatile <16 x i8> %a.add, ptr %a650 ret void651}652 653define void @generic_volatile_2xi16(ptr %a) {654; CHECK-LABEL: generic_volatile_2xi16(655; CHECK: {656; CHECK-NEXT: .reg .b16 %rs<5>;657; CHECK-NEXT: .reg .b64 %rd<2>;658; CHECK-EMPTY:659; CHECK-NEXT: // %bb.0:660; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xi16_param_0];661; CHECK-NEXT: ld.volatile.v2.b16 {%rs1, %rs2}, [%rd1];662; CHECK-NEXT: add.s16 %rs3, %rs2, 1;663; CHECK-NEXT: add.s16 %rs4, %rs1, 1;664; CHECK-NEXT: st.volatile.v2.b16 [%rd1], {%rs4, %rs3};665; CHECK-NEXT: ret;666 %a.load = load volatile <2 x i16>, ptr %a667 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>668 store volatile <2 x i16> %a.add, ptr %a669 ret void670}671 672define void @generic_volatile_4xi16(ptr %a) {673; CHECK-LABEL: generic_volatile_4xi16(674; CHECK: {675; CHECK-NEXT: .reg .b16 %rs<9>;676; CHECK-NEXT: .reg .b64 %rd<2>;677; CHECK-EMPTY:678; CHECK-NEXT: // %bb.0:679; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xi16_param_0];680; CHECK-NEXT: ld.volatile.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];681; CHECK-NEXT: add.s16 %rs5, %rs4, 1;682; CHECK-NEXT: add.s16 %rs6, %rs3, 1;683; CHECK-NEXT: add.s16 %rs7, %rs2, 1;684; CHECK-NEXT: add.s16 %rs8, %rs1, 1;685; CHECK-NEXT: st.volatile.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};686; CHECK-NEXT: ret;687 %a.load = load volatile <4 x i16>, ptr %a688 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>689 store volatile <4 x i16> %a.add, ptr %a690 ret void691}692 693define void @generic_volatile_8xi16(ptr %a) {694; CHECK-LABEL: generic_volatile_8xi16(695; CHECK: {696; CHECK-NEXT: .reg .b16 %rs<17>;697; CHECK-NEXT: .reg .b32 %r<9>;698; CHECK-NEXT: .reg .b64 %rd<2>;699; CHECK-EMPTY:700; CHECK-NEXT: // %bb.0:701; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_8xi16_param_0];702; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];703; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;704; CHECK-NEXT: add.s16 %rs3, %rs2, 1;705; CHECK-NEXT: add.s16 %rs4, %rs1, 1;706; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};707; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;708; CHECK-NEXT: add.s16 %rs7, %rs6, 1;709; CHECK-NEXT: add.s16 %rs8, %rs5, 1;710; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};711; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;712; CHECK-NEXT: add.s16 %rs11, %rs10, 1;713; CHECK-NEXT: add.s16 %rs12, %rs9, 1;714; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};715; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;716; CHECK-NEXT: add.s16 %rs15, %rs14, 1;717; CHECK-NEXT: add.s16 %rs16, %rs13, 1;718; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};719; CHECK-NEXT: st.volatile.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};720; CHECK-NEXT: ret;721 %a.load = load volatile <8 x i16>, ptr %a722 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>723 store volatile <8 x i16> %a.add, ptr %a724 ret void725}726 727define void @generic_volatile_2xi32(ptr %a) {728; CHECK-LABEL: generic_volatile_2xi32(729; CHECK: {730; CHECK-NEXT: .reg .b32 %r<5>;731; CHECK-NEXT: .reg .b64 %rd<2>;732; CHECK-EMPTY:733; CHECK-NEXT: // %bb.0:734; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xi32_param_0];735; CHECK-NEXT: ld.volatile.v2.b32 {%r1, %r2}, [%rd1];736; CHECK-NEXT: add.s32 %r3, %r2, 1;737; CHECK-NEXT: add.s32 %r4, %r1, 1;738; CHECK-NEXT: st.volatile.v2.b32 [%rd1], {%r4, %r3};739; CHECK-NEXT: ret;740 %a.load = load volatile <2 x i32>, ptr %a741 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>742 store volatile <2 x i32> %a.add, ptr %a743 ret void744}745 746define void @generic_volatile_4xi32(ptr %a) {747; CHECK-LABEL: generic_volatile_4xi32(748; CHECK: {749; CHECK-NEXT: .reg .b32 %r<9>;750; CHECK-NEXT: .reg .b64 %rd<2>;751; CHECK-EMPTY:752; CHECK-NEXT: // %bb.0:753; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xi32_param_0];754; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];755; CHECK-NEXT: add.s32 %r5, %r4, 1;756; CHECK-NEXT: add.s32 %r6, %r3, 1;757; CHECK-NEXT: add.s32 %r7, %r2, 1;758; CHECK-NEXT: add.s32 %r8, %r1, 1;759; CHECK-NEXT: st.volatile.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};760; CHECK-NEXT: ret;761 %a.load = load volatile <4 x i32>, ptr %a762 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>763 store volatile <4 x i32> %a.add, ptr %a764 ret void765}766 767define void @generic_volatile_2xi64(ptr %a) {768; CHECK-LABEL: generic_volatile_2xi64(769; CHECK: {770; CHECK-NEXT: .reg .b64 %rd<6>;771; CHECK-EMPTY:772; CHECK-NEXT: // %bb.0:773; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xi64_param_0];774; CHECK-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];775; CHECK-NEXT: add.s64 %rd4, %rd3, 1;776; CHECK-NEXT: add.s64 %rd5, %rd2, 1;777; CHECK-NEXT: st.volatile.v2.b64 [%rd1], {%rd5, %rd4};778; CHECK-NEXT: ret;779 %a.load = load volatile <2 x i64>, ptr %a780 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>781 store volatile <2 x i64> %a.add, ptr %a782 ret void783}784 785define void @generic_volatile_2xfloat(ptr %a) {786; CHECK-LABEL: generic_volatile_2xfloat(787; CHECK: {788; CHECK-NEXT: .reg .b32 %r<5>;789; CHECK-NEXT: .reg .b64 %rd<2>;790; CHECK-EMPTY:791; CHECK-NEXT: // %bb.0:792; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xfloat_param_0];793; CHECK-NEXT: ld.volatile.v2.b32 {%r1, %r2}, [%rd1];794; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;795; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;796; CHECK-NEXT: st.volatile.v2.b32 [%rd1], {%r4, %r3};797; CHECK-NEXT: ret;798 %a.load = load volatile <2 x float>, ptr %a799 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>800 store volatile <2 x float> %a.add, ptr %a801 ret void802}803 804define void @generic_volatile_4xfloat(ptr %a) {805; CHECK-LABEL: generic_volatile_4xfloat(806; CHECK: {807; CHECK-NEXT: .reg .b32 %r<9>;808; CHECK-NEXT: .reg .b64 %rd<2>;809; CHECK-EMPTY:810; CHECK-NEXT: // %bb.0:811; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xfloat_param_0];812; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];813; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;814; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;815; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;816; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;817; CHECK-NEXT: st.volatile.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};818; CHECK-NEXT: ret;819 %a.load = load volatile <4 x float>, ptr %a820 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>821 store volatile <4 x float> %a.add, ptr %a822 ret void823}824 825define void @generic_volatile_2xdouble(ptr %a) {826; CHECK-LABEL: generic_volatile_2xdouble(827; CHECK: {828; CHECK-NEXT: .reg .b64 %rd<6>;829; CHECK-EMPTY:830; CHECK-NEXT: // %bb.0:831; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_2xdouble_param_0];832; CHECK-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];833; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;834; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;835; CHECK-NEXT: st.volatile.v2.b64 [%rd1], {%rd5, %rd4};836; CHECK-NEXT: ret;837 %a.load = load volatile <2 x double>, ptr %a838 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>839 store volatile <2 x double> %a.add, ptr %a840 ret void841}842 843;; global statespace844 845; global846 847define void @global_2xi8(ptr addrspace(1) %a) {848; CHECK-LABEL: global_2xi8(849; CHECK: {850; CHECK-NEXT: .reg .b16 %rs<5>;851; CHECK-NEXT: .reg .b64 %rd<2>;852; CHECK-EMPTY:853; CHECK-NEXT: // %bb.0:854; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi8_param_0];855; CHECK-NEXT: ld.global.v2.b8 {%rs1, %rs2}, [%rd1];856; CHECK-NEXT: add.s16 %rs3, %rs2, 1;857; CHECK-NEXT: add.s16 %rs4, %rs1, 1;858; CHECK-NEXT: st.global.v2.b8 [%rd1], {%rs4, %rs3};859; CHECK-NEXT: ret;860 %a.load = load <2 x i8>, ptr addrspace(1) %a861 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>862 store <2 x i8> %a.add, ptr addrspace(1) %a863 ret void864}865 866define void @global_4xi8(ptr addrspace(1) %a) {867; CHECK-LABEL: global_4xi8(868; CHECK: {869; CHECK-NEXT: .reg .b16 %rs<9>;870; CHECK-NEXT: .reg .b32 %r<13>;871; CHECK-NEXT: .reg .b64 %rd<2>;872; CHECK-EMPTY:873; CHECK-NEXT: // %bb.0:874; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_param_0];875; CHECK-NEXT: ld.global.b32 %r1, [%rd1];876; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;877; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;878; CHECK-NEXT: add.s16 %rs2, %rs1, 1;879; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;880; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;881; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;882; CHECK-NEXT: add.s16 %rs4, %rs3, 1;883; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;884; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;885; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;886; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;887; CHECK-NEXT: add.s16 %rs6, %rs5, 1;888; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;889; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;890; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;891; CHECK-NEXT: add.s16 %rs8, %rs7, 1;892; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;893; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;894; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;895; CHECK-NEXT: st.global.b32 [%rd1], %r12;896; CHECK-NEXT: ret;897 %a.load = load <4 x i8>, ptr addrspace(1) %a898 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>899 store <4 x i8> %a.add, ptr addrspace(1) %a900 ret void901}902 903define void @global_8xi8(ptr addrspace(1) %a) {904; CHECK-LABEL: global_8xi8(905; CHECK: {906; CHECK-NEXT: .reg .b16 %rs<17>;907; CHECK-NEXT: .reg .b32 %r<25>;908; CHECK-NEXT: .reg .b64 %rd<2>;909; CHECK-EMPTY:910; CHECK-NEXT: // %bb.0:911; CHECK-NEXT: ld.param.b64 %rd1, [global_8xi8_param_0];912; CHECK-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd1];913; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;914; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;915; CHECK-NEXT: add.s16 %rs2, %rs1, 1;916; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;917; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;918; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;919; CHECK-NEXT: add.s16 %rs4, %rs3, 1;920; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;921; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;922; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;923; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;924; CHECK-NEXT: add.s16 %rs6, %rs5, 1;925; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;926; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;927; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;928; CHECK-NEXT: add.s16 %rs8, %rs7, 1;929; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;930; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;931; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;932; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;933; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;934; CHECK-NEXT: add.s16 %rs10, %rs9, 1;935; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;936; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;937; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;938; CHECK-NEXT: add.s16 %rs12, %rs11, 1;939; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;940; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;941; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;942; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;943; CHECK-NEXT: add.s16 %rs14, %rs13, 1;944; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;945; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;946; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;947; CHECK-NEXT: add.s16 %rs16, %rs15, 1;948; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;949; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;950; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;951; CHECK-NEXT: st.global.v2.b32 [%rd1], {%r24, %r13};952; CHECK-NEXT: ret;953 %a.load = load <8 x i8>, ptr addrspace(1) %a954 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>955 store <8 x i8> %a.add, ptr addrspace(1) %a956 ret void957}958 959define void @global_16xi8(ptr addrspace(1) %a) {960; CHECK-LABEL: global_16xi8(961; CHECK: {962; CHECK-NEXT: .reg .b16 %rs<33>;963; CHECK-NEXT: .reg .b32 %r<49>;964; CHECK-NEXT: .reg .b64 %rd<2>;965; CHECK-EMPTY:966; CHECK-NEXT: // %bb.0:967; CHECK-NEXT: ld.param.b64 %rd1, [global_16xi8_param_0];968; CHECK-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];969; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;970; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;971; CHECK-NEXT: add.s16 %rs2, %rs1, 1;972; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;973; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;974; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;975; CHECK-NEXT: add.s16 %rs4, %rs3, 1;976; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;977; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;978; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;979; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;980; CHECK-NEXT: add.s16 %rs6, %rs5, 1;981; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;982; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;983; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;984; CHECK-NEXT: add.s16 %rs8, %rs7, 1;985; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;986; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;987; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;988; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;989; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;990; CHECK-NEXT: add.s16 %rs10, %rs9, 1;991; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;992; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;993; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;994; CHECK-NEXT: add.s16 %rs12, %rs11, 1;995; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;996; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;997; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;998; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;999; CHECK-NEXT: add.s16 %rs14, %rs13, 1;1000; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;1001; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;1002; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;1003; CHECK-NEXT: add.s16 %rs16, %rs15, 1;1004; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;1005; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;1006; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;1007; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;1008; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;1009; CHECK-NEXT: add.s16 %rs18, %rs17, 1;1010; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;1011; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;1012; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;1013; CHECK-NEXT: add.s16 %rs20, %rs19, 1;1014; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;1015; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;1016; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;1017; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;1018; CHECK-NEXT: add.s16 %rs22, %rs21, 1;1019; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;1020; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;1021; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;1022; CHECK-NEXT: add.s16 %rs24, %rs23, 1;1023; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;1024; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;1025; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;1026; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;1027; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;1028; CHECK-NEXT: add.s16 %rs26, %rs25, 1;1029; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;1030; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;1031; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;1032; CHECK-NEXT: add.s16 %rs28, %rs27, 1;1033; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;1034; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;1035; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;1036; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;1037; CHECK-NEXT: add.s16 %rs30, %rs29, 1;1038; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;1039; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;1040; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;1041; CHECK-NEXT: add.s16 %rs32, %rs31, 1;1042; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;1043; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;1044; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;1045; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};1046; CHECK-NEXT: ret;1047 %a.load = load <16 x i8>, ptr addrspace(1) %a1048 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>1049 store <16 x i8> %a.add, ptr addrspace(1) %a1050 ret void1051}1052 1053define void @global_2xi16(ptr addrspace(1) %a) {1054; CHECK-LABEL: global_2xi16(1055; CHECK: {1056; CHECK-NEXT: .reg .b16 %rs<5>;1057; CHECK-NEXT: .reg .b64 %rd<2>;1058; CHECK-EMPTY:1059; CHECK-NEXT: // %bb.0:1060; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_param_0];1061; CHECK-NEXT: ld.global.v2.b16 {%rs1, %rs2}, [%rd1];1062; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1063; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1064; CHECK-NEXT: st.global.v2.b16 [%rd1], {%rs4, %rs3};1065; CHECK-NEXT: ret;1066 %a.load = load <2 x i16>, ptr addrspace(1) %a1067 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>1068 store <2 x i16> %a.add, ptr addrspace(1) %a1069 ret void1070}1071 1072define void @global_4xi16(ptr addrspace(1) %a) {1073; CHECK-LABEL: global_4xi16(1074; CHECK: {1075; CHECK-NEXT: .reg .b16 %rs<9>;1076; CHECK-NEXT: .reg .b64 %rd<2>;1077; CHECK-EMPTY:1078; CHECK-NEXT: // %bb.0:1079; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi16_param_0];1080; CHECK-NEXT: ld.global.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];1081; CHECK-NEXT: add.s16 %rs5, %rs4, 1;1082; CHECK-NEXT: add.s16 %rs6, %rs3, 1;1083; CHECK-NEXT: add.s16 %rs7, %rs2, 1;1084; CHECK-NEXT: add.s16 %rs8, %rs1, 1;1085; CHECK-NEXT: st.global.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};1086; CHECK-NEXT: ret;1087 %a.load = load <4 x i16>, ptr addrspace(1) %a1088 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>1089 store <4 x i16> %a.add, ptr addrspace(1) %a1090 ret void1091}1092 1093define void @global_8xi16(ptr addrspace(1) %a) {1094; CHECK-LABEL: global_8xi16(1095; CHECK: {1096; CHECK-NEXT: .reg .b16 %rs<17>;1097; CHECK-NEXT: .reg .b32 %r<9>;1098; CHECK-NEXT: .reg .b64 %rd<2>;1099; CHECK-EMPTY:1100; CHECK-NEXT: // %bb.0:1101; CHECK-NEXT: ld.param.b64 %rd1, [global_8xi16_param_0];1102; CHECK-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1103; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;1104; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1105; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1106; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};1107; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;1108; CHECK-NEXT: add.s16 %rs7, %rs6, 1;1109; CHECK-NEXT: add.s16 %rs8, %rs5, 1;1110; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};1111; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;1112; CHECK-NEXT: add.s16 %rs11, %rs10, 1;1113; CHECK-NEXT: add.s16 %rs12, %rs9, 1;1114; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};1115; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;1116; CHECK-NEXT: add.s16 %rs15, %rs14, 1;1117; CHECK-NEXT: add.s16 %rs16, %rs13, 1;1118; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};1119; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1120; CHECK-NEXT: ret;1121 %a.load = load <8 x i16>, ptr addrspace(1) %a1122 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>1123 store <8 x i16> %a.add, ptr addrspace(1) %a1124 ret void1125}1126 1127define void @global_2xi32(ptr addrspace(1) %a) {1128; CHECK-LABEL: global_2xi32(1129; CHECK: {1130; CHECK-NEXT: .reg .b32 %r<5>;1131; CHECK-NEXT: .reg .b64 %rd<2>;1132; CHECK-EMPTY:1133; CHECK-NEXT: // %bb.0:1134; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi32_param_0];1135; CHECK-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd1];1136; CHECK-NEXT: add.s32 %r3, %r2, 1;1137; CHECK-NEXT: add.s32 %r4, %r1, 1;1138; CHECK-NEXT: st.global.v2.b32 [%rd1], {%r4, %r3};1139; CHECK-NEXT: ret;1140 %a.load = load <2 x i32>, ptr addrspace(1) %a1141 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>1142 store <2 x i32> %a.add, ptr addrspace(1) %a1143 ret void1144}1145 1146define void @global_4xi32(ptr addrspace(1) %a) {1147; CHECK-LABEL: global_4xi32(1148; CHECK: {1149; CHECK-NEXT: .reg .b32 %r<9>;1150; CHECK-NEXT: .reg .b64 %rd<2>;1151; CHECK-EMPTY:1152; CHECK-NEXT: // %bb.0:1153; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi32_param_0];1154; CHECK-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1155; CHECK-NEXT: add.s32 %r5, %r4, 1;1156; CHECK-NEXT: add.s32 %r6, %r3, 1;1157; CHECK-NEXT: add.s32 %r7, %r2, 1;1158; CHECK-NEXT: add.s32 %r8, %r1, 1;1159; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1160; CHECK-NEXT: ret;1161 %a.load = load <4 x i32>, ptr addrspace(1) %a1162 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>1163 store <4 x i32> %a.add, ptr addrspace(1) %a1164 ret void1165}1166 1167define void @global_2xi64(ptr addrspace(1) %a) {1168; CHECK-LABEL: global_2xi64(1169; CHECK: {1170; CHECK-NEXT: .reg .b64 %rd<6>;1171; CHECK-EMPTY:1172; CHECK-NEXT: // %bb.0:1173; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi64_param_0];1174; CHECK-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];1175; CHECK-NEXT: add.s64 %rd4, %rd3, 1;1176; CHECK-NEXT: add.s64 %rd5, %rd2, 1;1177; CHECK-NEXT: st.global.v2.b64 [%rd1], {%rd5, %rd4};1178; CHECK-NEXT: ret;1179 %a.load = load <2 x i64>, ptr addrspace(1) %a1180 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>1181 store <2 x i64> %a.add, ptr addrspace(1) %a1182 ret void1183}1184 1185define void @global_2xfloat(ptr addrspace(1) %a) {1186; CHECK-LABEL: global_2xfloat(1187; CHECK: {1188; CHECK-NEXT: .reg .b32 %r<5>;1189; CHECK-NEXT: .reg .b64 %rd<2>;1190; CHECK-EMPTY:1191; CHECK-NEXT: // %bb.0:1192; CHECK-NEXT: ld.param.b64 %rd1, [global_2xfloat_param_0];1193; CHECK-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd1];1194; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;1195; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;1196; CHECK-NEXT: st.global.v2.b32 [%rd1], {%r4, %r3};1197; CHECK-NEXT: ret;1198 %a.load = load <2 x float>, ptr addrspace(1) %a1199 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>1200 store <2 x float> %a.add, ptr addrspace(1) %a1201 ret void1202}1203 1204define void @global_4xfloat(ptr addrspace(1) %a) {1205; CHECK-LABEL: global_4xfloat(1206; CHECK: {1207; CHECK-NEXT: .reg .b32 %r<9>;1208; CHECK-NEXT: .reg .b64 %rd<2>;1209; CHECK-EMPTY:1210; CHECK-NEXT: // %bb.0:1211; CHECK-NEXT: ld.param.b64 %rd1, [global_4xfloat_param_0];1212; CHECK-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1213; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;1214; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;1215; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;1216; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;1217; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1218; CHECK-NEXT: ret;1219 %a.load = load <4 x float>, ptr addrspace(1) %a1220 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>1221 store <4 x float> %a.add, ptr addrspace(1) %a1222 ret void1223}1224 1225define void @global_2xdouble(ptr addrspace(1) %a) {1226; CHECK-LABEL: global_2xdouble(1227; CHECK: {1228; CHECK-NEXT: .reg .b64 %rd<6>;1229; CHECK-EMPTY:1230; CHECK-NEXT: // %bb.0:1231; CHECK-NEXT: ld.param.b64 %rd1, [global_2xdouble_param_0];1232; CHECK-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];1233; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;1234; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;1235; CHECK-NEXT: st.global.v2.b64 [%rd1], {%rd5, %rd4};1236; CHECK-NEXT: ret;1237 %a.load = load <2 x double>, ptr addrspace(1) %a1238 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>1239 store <2 x double> %a.add, ptr addrspace(1) %a1240 ret void1241}1242 1243; global_volatile1244 1245define void @global_volatile_2xi8(ptr addrspace(1) %a) {1246; CHECK-LABEL: global_volatile_2xi8(1247; CHECK: {1248; CHECK-NEXT: .reg .b16 %rs<5>;1249; CHECK-NEXT: .reg .b64 %rd<2>;1250; CHECK-EMPTY:1251; CHECK-NEXT: // %bb.0:1252; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xi8_param_0];1253; CHECK-NEXT: ld.volatile.global.v2.b8 {%rs1, %rs2}, [%rd1];1254; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1255; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1256; CHECK-NEXT: st.volatile.global.v2.b8 [%rd1], {%rs4, %rs3};1257; CHECK-NEXT: ret;1258 %a.load = load volatile <2 x i8>, ptr addrspace(1) %a1259 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>1260 store volatile <2 x i8> %a.add, ptr addrspace(1) %a1261 ret void1262}1263 1264define void @global_volatile_4xi8(ptr addrspace(1) %a) {1265; CHECK-LABEL: global_volatile_4xi8(1266; CHECK: {1267; CHECK-NEXT: .reg .b16 %rs<9>;1268; CHECK-NEXT: .reg .b32 %r<13>;1269; CHECK-NEXT: .reg .b64 %rd<2>;1270; CHECK-EMPTY:1271; CHECK-NEXT: // %bb.0:1272; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_4xi8_param_0];1273; CHECK-NEXT: ld.volatile.global.b32 %r1, [%rd1];1274; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;1275; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;1276; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1277; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;1278; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;1279; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;1280; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1281; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;1282; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;1283; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;1284; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;1285; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1286; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;1287; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;1288; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;1289; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1290; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;1291; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;1292; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;1293; CHECK-NEXT: st.volatile.global.b32 [%rd1], %r12;1294; CHECK-NEXT: ret;1295 %a.load = load volatile <4 x i8>, ptr addrspace(1) %a1296 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>1297 store volatile <4 x i8> %a.add, ptr addrspace(1) %a1298 ret void1299}1300 1301define void @global_volatile_8xi8(ptr addrspace(1) %a) {1302; CHECK-LABEL: global_volatile_8xi8(1303; CHECK: {1304; CHECK-NEXT: .reg .b16 %rs<17>;1305; CHECK-NEXT: .reg .b32 %r<25>;1306; CHECK-NEXT: .reg .b64 %rd<2>;1307; CHECK-EMPTY:1308; CHECK-NEXT: // %bb.0:1309; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_8xi8_param_0];1310; CHECK-NEXT: ld.volatile.global.v2.b32 {%r1, %r2}, [%rd1];1311; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;1312; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;1313; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1314; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;1315; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;1316; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;1317; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1318; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;1319; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;1320; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;1321; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;1322; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1323; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;1324; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;1325; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;1326; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1327; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;1328; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;1329; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;1330; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;1331; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;1332; CHECK-NEXT: add.s16 %rs10, %rs9, 1;1333; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;1334; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;1335; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;1336; CHECK-NEXT: add.s16 %rs12, %rs11, 1;1337; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;1338; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;1339; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;1340; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;1341; CHECK-NEXT: add.s16 %rs14, %rs13, 1;1342; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;1343; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;1344; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;1345; CHECK-NEXT: add.s16 %rs16, %rs15, 1;1346; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;1347; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;1348; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;1349; CHECK-NEXT: st.volatile.global.v2.b32 [%rd1], {%r24, %r13};1350; CHECK-NEXT: ret;1351 %a.load = load volatile <8 x i8>, ptr addrspace(1) %a1352 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>1353 store volatile <8 x i8> %a.add, ptr addrspace(1) %a1354 ret void1355}1356 1357define void @global_volatile_16xi8(ptr addrspace(1) %a) {1358; CHECK-LABEL: global_volatile_16xi8(1359; CHECK: {1360; CHECK-NEXT: .reg .b16 %rs<33>;1361; CHECK-NEXT: .reg .b32 %r<49>;1362; CHECK-NEXT: .reg .b64 %rd<2>;1363; CHECK-EMPTY:1364; CHECK-NEXT: // %bb.0:1365; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_16xi8_param_0];1366; CHECK-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1367; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;1368; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;1369; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1370; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;1371; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;1372; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;1373; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1374; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;1375; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;1376; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;1377; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;1378; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1379; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;1380; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;1381; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;1382; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1383; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;1384; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;1385; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;1386; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;1387; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;1388; CHECK-NEXT: add.s16 %rs10, %rs9, 1;1389; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;1390; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;1391; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;1392; CHECK-NEXT: add.s16 %rs12, %rs11, 1;1393; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;1394; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;1395; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;1396; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;1397; CHECK-NEXT: add.s16 %rs14, %rs13, 1;1398; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;1399; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;1400; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;1401; CHECK-NEXT: add.s16 %rs16, %rs15, 1;1402; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;1403; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;1404; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;1405; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;1406; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;1407; CHECK-NEXT: add.s16 %rs18, %rs17, 1;1408; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;1409; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;1410; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;1411; CHECK-NEXT: add.s16 %rs20, %rs19, 1;1412; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;1413; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;1414; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;1415; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;1416; CHECK-NEXT: add.s16 %rs22, %rs21, 1;1417; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;1418; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;1419; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;1420; CHECK-NEXT: add.s16 %rs24, %rs23, 1;1421; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;1422; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;1423; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;1424; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;1425; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;1426; CHECK-NEXT: add.s16 %rs26, %rs25, 1;1427; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;1428; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;1429; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;1430; CHECK-NEXT: add.s16 %rs28, %rs27, 1;1431; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;1432; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;1433; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;1434; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;1435; CHECK-NEXT: add.s16 %rs30, %rs29, 1;1436; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;1437; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;1438; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;1439; CHECK-NEXT: add.s16 %rs32, %rs31, 1;1440; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;1441; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;1442; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;1443; CHECK-NEXT: st.volatile.global.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};1444; CHECK-NEXT: ret;1445 %a.load = load volatile <16 x i8>, ptr addrspace(1) %a1446 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>1447 store volatile <16 x i8> %a.add, ptr addrspace(1) %a1448 ret void1449}1450 1451define void @global_volatile_2xi16(ptr addrspace(1) %a) {1452; CHECK-LABEL: global_volatile_2xi16(1453; CHECK: {1454; CHECK-NEXT: .reg .b16 %rs<5>;1455; CHECK-NEXT: .reg .b64 %rd<2>;1456; CHECK-EMPTY:1457; CHECK-NEXT: // %bb.0:1458; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xi16_param_0];1459; CHECK-NEXT: ld.volatile.global.v2.b16 {%rs1, %rs2}, [%rd1];1460; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1461; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1462; CHECK-NEXT: st.volatile.global.v2.b16 [%rd1], {%rs4, %rs3};1463; CHECK-NEXT: ret;1464 %a.load = load volatile <2 x i16>, ptr addrspace(1) %a1465 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>1466 store volatile <2 x i16> %a.add, ptr addrspace(1) %a1467 ret void1468}1469 1470define void @global_volatile_4xi16(ptr addrspace(1) %a) {1471; CHECK-LABEL: global_volatile_4xi16(1472; CHECK: {1473; CHECK-NEXT: .reg .b16 %rs<9>;1474; CHECK-NEXT: .reg .b64 %rd<2>;1475; CHECK-EMPTY:1476; CHECK-NEXT: // %bb.0:1477; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_4xi16_param_0];1478; CHECK-NEXT: ld.volatile.global.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];1479; CHECK-NEXT: add.s16 %rs5, %rs4, 1;1480; CHECK-NEXT: add.s16 %rs6, %rs3, 1;1481; CHECK-NEXT: add.s16 %rs7, %rs2, 1;1482; CHECK-NEXT: add.s16 %rs8, %rs1, 1;1483; CHECK-NEXT: st.volatile.global.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};1484; CHECK-NEXT: ret;1485 %a.load = load volatile <4 x i16>, ptr addrspace(1) %a1486 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>1487 store volatile <4 x i16> %a.add, ptr addrspace(1) %a1488 ret void1489}1490 1491define void @global_volatile_8xi16(ptr addrspace(1) %a) {1492; CHECK-LABEL: global_volatile_8xi16(1493; CHECK: {1494; CHECK-NEXT: .reg .b16 %rs<17>;1495; CHECK-NEXT: .reg .b32 %r<9>;1496; CHECK-NEXT: .reg .b64 %rd<2>;1497; CHECK-EMPTY:1498; CHECK-NEXT: // %bb.0:1499; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_8xi16_param_0];1500; CHECK-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1501; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;1502; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1503; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1504; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};1505; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;1506; CHECK-NEXT: add.s16 %rs7, %rs6, 1;1507; CHECK-NEXT: add.s16 %rs8, %rs5, 1;1508; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};1509; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;1510; CHECK-NEXT: add.s16 %rs11, %rs10, 1;1511; CHECK-NEXT: add.s16 %rs12, %rs9, 1;1512; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};1513; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;1514; CHECK-NEXT: add.s16 %rs15, %rs14, 1;1515; CHECK-NEXT: add.s16 %rs16, %rs13, 1;1516; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};1517; CHECK-NEXT: st.volatile.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1518; CHECK-NEXT: ret;1519 %a.load = load volatile <8 x i16>, ptr addrspace(1) %a1520 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>1521 store volatile <8 x i16> %a.add, ptr addrspace(1) %a1522 ret void1523}1524 1525define void @global_volatile_2xi32(ptr addrspace(1) %a) {1526; CHECK-LABEL: global_volatile_2xi32(1527; CHECK: {1528; CHECK-NEXT: .reg .b32 %r<5>;1529; CHECK-NEXT: .reg .b64 %rd<2>;1530; CHECK-EMPTY:1531; CHECK-NEXT: // %bb.0:1532; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xi32_param_0];1533; CHECK-NEXT: ld.volatile.global.v2.b32 {%r1, %r2}, [%rd1];1534; CHECK-NEXT: add.s32 %r3, %r2, 1;1535; CHECK-NEXT: add.s32 %r4, %r1, 1;1536; CHECK-NEXT: st.volatile.global.v2.b32 [%rd1], {%r4, %r3};1537; CHECK-NEXT: ret;1538 %a.load = load volatile <2 x i32>, ptr addrspace(1) %a1539 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>1540 store volatile <2 x i32> %a.add, ptr addrspace(1) %a1541 ret void1542}1543 1544define void @global_volatile_4xi32(ptr addrspace(1) %a) {1545; CHECK-LABEL: global_volatile_4xi32(1546; CHECK: {1547; CHECK-NEXT: .reg .b32 %r<9>;1548; CHECK-NEXT: .reg .b64 %rd<2>;1549; CHECK-EMPTY:1550; CHECK-NEXT: // %bb.0:1551; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_4xi32_param_0];1552; CHECK-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1553; CHECK-NEXT: add.s32 %r5, %r4, 1;1554; CHECK-NEXT: add.s32 %r6, %r3, 1;1555; CHECK-NEXT: add.s32 %r7, %r2, 1;1556; CHECK-NEXT: add.s32 %r8, %r1, 1;1557; CHECK-NEXT: st.volatile.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1558; CHECK-NEXT: ret;1559 %a.load = load volatile <4 x i32>, ptr addrspace(1) %a1560 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>1561 store volatile <4 x i32> %a.add, ptr addrspace(1) %a1562 ret void1563}1564 1565define void @global_volatile_2xi64(ptr addrspace(1) %a) {1566; CHECK-LABEL: global_volatile_2xi64(1567; CHECK: {1568; CHECK-NEXT: .reg .b64 %rd<6>;1569; CHECK-EMPTY:1570; CHECK-NEXT: // %bb.0:1571; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xi64_param_0];1572; CHECK-NEXT: ld.volatile.global.v2.b64 {%rd2, %rd3}, [%rd1];1573; CHECK-NEXT: add.s64 %rd4, %rd3, 1;1574; CHECK-NEXT: add.s64 %rd5, %rd2, 1;1575; CHECK-NEXT: st.volatile.global.v2.b64 [%rd1], {%rd5, %rd4};1576; CHECK-NEXT: ret;1577 %a.load = load volatile <2 x i64>, ptr addrspace(1) %a1578 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>1579 store volatile <2 x i64> %a.add, ptr addrspace(1) %a1580 ret void1581}1582 1583define void @global_volatile_2xfloat(ptr addrspace(1) %a) {1584; CHECK-LABEL: global_volatile_2xfloat(1585; CHECK: {1586; CHECK-NEXT: .reg .b32 %r<5>;1587; CHECK-NEXT: .reg .b64 %rd<2>;1588; CHECK-EMPTY:1589; CHECK-NEXT: // %bb.0:1590; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xfloat_param_0];1591; CHECK-NEXT: ld.volatile.global.v2.b32 {%r1, %r2}, [%rd1];1592; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;1593; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;1594; CHECK-NEXT: st.volatile.global.v2.b32 [%rd1], {%r4, %r3};1595; CHECK-NEXT: ret;1596 %a.load = load volatile <2 x float>, ptr addrspace(1) %a1597 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>1598 store volatile <2 x float> %a.add, ptr addrspace(1) %a1599 ret void1600}1601 1602define void @global_volatile_4xfloat(ptr addrspace(1) %a) {1603; CHECK-LABEL: global_volatile_4xfloat(1604; CHECK: {1605; CHECK-NEXT: .reg .b32 %r<9>;1606; CHECK-NEXT: .reg .b64 %rd<2>;1607; CHECK-EMPTY:1608; CHECK-NEXT: // %bb.0:1609; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_4xfloat_param_0];1610; CHECK-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1611; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;1612; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;1613; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;1614; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;1615; CHECK-NEXT: st.volatile.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1616; CHECK-NEXT: ret;1617 %a.load = load volatile <4 x float>, ptr addrspace(1) %a1618 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>1619 store volatile <4 x float> %a.add, ptr addrspace(1) %a1620 ret void1621}1622 1623define void @global_volatile_2xdouble(ptr addrspace(1) %a) {1624; CHECK-LABEL: global_volatile_2xdouble(1625; CHECK: {1626; CHECK-NEXT: .reg .b64 %rd<6>;1627; CHECK-EMPTY:1628; CHECK-NEXT: // %bb.0:1629; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_2xdouble_param_0];1630; CHECK-NEXT: ld.volatile.global.v2.b64 {%rd2, %rd3}, [%rd1];1631; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;1632; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;1633; CHECK-NEXT: st.volatile.global.v2.b64 [%rd1], {%rd5, %rd4};1634; CHECK-NEXT: ret;1635 %a.load = load volatile <2 x double>, ptr addrspace(1) %a1636 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>1637 store volatile <2 x double> %a.add, ptr addrspace(1) %a1638 ret void1639}1640 1641;; shared statespace1642 1643; shared1644 1645define void @shared_2xi8(ptr addrspace(3) %a) {1646; CHECK-LABEL: shared_2xi8(1647; CHECK: {1648; CHECK-NEXT: .reg .b16 %rs<5>;1649; CHECK-NEXT: .reg .b64 %rd<2>;1650; CHECK-EMPTY:1651; CHECK-NEXT: // %bb.0:1652; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xi8_param_0];1653; CHECK-NEXT: ld.shared.v2.b8 {%rs1, %rs2}, [%rd1];1654; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1655; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1656; CHECK-NEXT: st.shared.v2.b8 [%rd1], {%rs4, %rs3};1657; CHECK-NEXT: ret;1658 %a.load = load <2 x i8>, ptr addrspace(3) %a1659 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>1660 store <2 x i8> %a.add, ptr addrspace(3) %a1661 ret void1662}1663 1664define void @shared_4xi8(ptr addrspace(3) %a) {1665; CHECK-LABEL: shared_4xi8(1666; CHECK: {1667; CHECK-NEXT: .reg .b16 %rs<9>;1668; CHECK-NEXT: .reg .b32 %r<13>;1669; CHECK-NEXT: .reg .b64 %rd<2>;1670; CHECK-EMPTY:1671; CHECK-NEXT: // %bb.0:1672; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xi8_param_0];1673; CHECK-NEXT: ld.shared.b32 %r1, [%rd1];1674; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;1675; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;1676; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1677; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;1678; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;1679; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;1680; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1681; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;1682; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;1683; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;1684; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;1685; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1686; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;1687; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;1688; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;1689; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1690; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;1691; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;1692; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;1693; CHECK-NEXT: st.shared.b32 [%rd1], %r12;1694; CHECK-NEXT: ret;1695 %a.load = load <4 x i8>, ptr addrspace(3) %a1696 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>1697 store <4 x i8> %a.add, ptr addrspace(3) %a1698 ret void1699}1700 1701define void @shared_8xi8(ptr addrspace(3) %a) {1702; CHECK-LABEL: shared_8xi8(1703; CHECK: {1704; CHECK-NEXT: .reg .b16 %rs<17>;1705; CHECK-NEXT: .reg .b32 %r<25>;1706; CHECK-NEXT: .reg .b64 %rd<2>;1707; CHECK-EMPTY:1708; CHECK-NEXT: // %bb.0:1709; CHECK-NEXT: ld.param.b64 %rd1, [shared_8xi8_param_0];1710; CHECK-NEXT: ld.shared.v2.b32 {%r1, %r2}, [%rd1];1711; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;1712; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;1713; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1714; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;1715; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;1716; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;1717; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1718; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;1719; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;1720; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;1721; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;1722; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1723; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;1724; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;1725; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;1726; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1727; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;1728; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;1729; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;1730; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;1731; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;1732; CHECK-NEXT: add.s16 %rs10, %rs9, 1;1733; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;1734; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;1735; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;1736; CHECK-NEXT: add.s16 %rs12, %rs11, 1;1737; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;1738; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;1739; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;1740; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;1741; CHECK-NEXT: add.s16 %rs14, %rs13, 1;1742; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;1743; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;1744; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;1745; CHECK-NEXT: add.s16 %rs16, %rs15, 1;1746; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;1747; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;1748; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;1749; CHECK-NEXT: st.shared.v2.b32 [%rd1], {%r24, %r13};1750; CHECK-NEXT: ret;1751 %a.load = load <8 x i8>, ptr addrspace(3) %a1752 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>1753 store <8 x i8> %a.add, ptr addrspace(3) %a1754 ret void1755}1756 1757define void @shared_16xi8(ptr addrspace(3) %a) {1758; CHECK-LABEL: shared_16xi8(1759; CHECK: {1760; CHECK-NEXT: .reg .b16 %rs<33>;1761; CHECK-NEXT: .reg .b32 %r<49>;1762; CHECK-NEXT: .reg .b64 %rd<2>;1763; CHECK-EMPTY:1764; CHECK-NEXT: // %bb.0:1765; CHECK-NEXT: ld.param.b64 %rd1, [shared_16xi8_param_0];1766; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1767; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;1768; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;1769; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1770; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;1771; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;1772; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;1773; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1774; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;1775; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;1776; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;1777; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;1778; CHECK-NEXT: add.s16 %rs6, %rs5, 1;1779; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;1780; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;1781; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;1782; CHECK-NEXT: add.s16 %rs8, %rs7, 1;1783; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;1784; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;1785; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;1786; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;1787; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;1788; CHECK-NEXT: add.s16 %rs10, %rs9, 1;1789; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;1790; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;1791; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;1792; CHECK-NEXT: add.s16 %rs12, %rs11, 1;1793; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;1794; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;1795; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;1796; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;1797; CHECK-NEXT: add.s16 %rs14, %rs13, 1;1798; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;1799; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;1800; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;1801; CHECK-NEXT: add.s16 %rs16, %rs15, 1;1802; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;1803; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;1804; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;1805; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;1806; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;1807; CHECK-NEXT: add.s16 %rs18, %rs17, 1;1808; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;1809; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;1810; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;1811; CHECK-NEXT: add.s16 %rs20, %rs19, 1;1812; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;1813; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;1814; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;1815; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;1816; CHECK-NEXT: add.s16 %rs22, %rs21, 1;1817; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;1818; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;1819; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;1820; CHECK-NEXT: add.s16 %rs24, %rs23, 1;1821; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;1822; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;1823; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;1824; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;1825; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;1826; CHECK-NEXT: add.s16 %rs26, %rs25, 1;1827; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;1828; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;1829; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;1830; CHECK-NEXT: add.s16 %rs28, %rs27, 1;1831; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;1832; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;1833; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;1834; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;1835; CHECK-NEXT: add.s16 %rs30, %rs29, 1;1836; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;1837; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;1838; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;1839; CHECK-NEXT: add.s16 %rs32, %rs31, 1;1840; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;1841; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;1842; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;1843; CHECK-NEXT: st.shared.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};1844; CHECK-NEXT: ret;1845 %a.load = load <16 x i8>, ptr addrspace(3) %a1846 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>1847 store <16 x i8> %a.add, ptr addrspace(3) %a1848 ret void1849}1850 1851define void @shared_2xi16(ptr addrspace(3) %a) {1852; CHECK-LABEL: shared_2xi16(1853; CHECK: {1854; CHECK-NEXT: .reg .b16 %rs<5>;1855; CHECK-NEXT: .reg .b64 %rd<2>;1856; CHECK-EMPTY:1857; CHECK-NEXT: // %bb.0:1858; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xi16_param_0];1859; CHECK-NEXT: ld.shared.v2.b16 {%rs1, %rs2}, [%rd1];1860; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1861; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1862; CHECK-NEXT: st.shared.v2.b16 [%rd1], {%rs4, %rs3};1863; CHECK-NEXT: ret;1864 %a.load = load <2 x i16>, ptr addrspace(3) %a1865 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>1866 store <2 x i16> %a.add, ptr addrspace(3) %a1867 ret void1868}1869 1870define void @shared_4xi16(ptr addrspace(3) %a) {1871; CHECK-LABEL: shared_4xi16(1872; CHECK: {1873; CHECK-NEXT: .reg .b16 %rs<9>;1874; CHECK-NEXT: .reg .b64 %rd<2>;1875; CHECK-EMPTY:1876; CHECK-NEXT: // %bb.0:1877; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xi16_param_0];1878; CHECK-NEXT: ld.shared.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];1879; CHECK-NEXT: add.s16 %rs5, %rs4, 1;1880; CHECK-NEXT: add.s16 %rs6, %rs3, 1;1881; CHECK-NEXT: add.s16 %rs7, %rs2, 1;1882; CHECK-NEXT: add.s16 %rs8, %rs1, 1;1883; CHECK-NEXT: st.shared.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};1884; CHECK-NEXT: ret;1885 %a.load = load <4 x i16>, ptr addrspace(3) %a1886 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>1887 store <4 x i16> %a.add, ptr addrspace(3) %a1888 ret void1889}1890 1891define void @shared_8xi16(ptr addrspace(3) %a) {1892; CHECK-LABEL: shared_8xi16(1893; CHECK: {1894; CHECK-NEXT: .reg .b16 %rs<17>;1895; CHECK-NEXT: .reg .b32 %r<9>;1896; CHECK-NEXT: .reg .b64 %rd<2>;1897; CHECK-EMPTY:1898; CHECK-NEXT: // %bb.0:1899; CHECK-NEXT: ld.param.b64 %rd1, [shared_8xi16_param_0];1900; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1901; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;1902; CHECK-NEXT: add.s16 %rs3, %rs2, 1;1903; CHECK-NEXT: add.s16 %rs4, %rs1, 1;1904; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};1905; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;1906; CHECK-NEXT: add.s16 %rs7, %rs6, 1;1907; CHECK-NEXT: add.s16 %rs8, %rs5, 1;1908; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};1909; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;1910; CHECK-NEXT: add.s16 %rs11, %rs10, 1;1911; CHECK-NEXT: add.s16 %rs12, %rs9, 1;1912; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};1913; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;1914; CHECK-NEXT: add.s16 %rs15, %rs14, 1;1915; CHECK-NEXT: add.s16 %rs16, %rs13, 1;1916; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};1917; CHECK-NEXT: st.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1918; CHECK-NEXT: ret;1919 %a.load = load <8 x i16>, ptr addrspace(3) %a1920 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>1921 store <8 x i16> %a.add, ptr addrspace(3) %a1922 ret void1923}1924 1925define void @shared_2xi32(ptr addrspace(3) %a) {1926; CHECK-LABEL: shared_2xi32(1927; CHECK: {1928; CHECK-NEXT: .reg .b32 %r<5>;1929; CHECK-NEXT: .reg .b64 %rd<2>;1930; CHECK-EMPTY:1931; CHECK-NEXT: // %bb.0:1932; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xi32_param_0];1933; CHECK-NEXT: ld.shared.v2.b32 {%r1, %r2}, [%rd1];1934; CHECK-NEXT: add.s32 %r3, %r2, 1;1935; CHECK-NEXT: add.s32 %r4, %r1, 1;1936; CHECK-NEXT: st.shared.v2.b32 [%rd1], {%r4, %r3};1937; CHECK-NEXT: ret;1938 %a.load = load <2 x i32>, ptr addrspace(3) %a1939 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>1940 store <2 x i32> %a.add, ptr addrspace(3) %a1941 ret void1942}1943 1944define void @shared_4xi32(ptr addrspace(3) %a) {1945; CHECK-LABEL: shared_4xi32(1946; CHECK: {1947; CHECK-NEXT: .reg .b32 %r<9>;1948; CHECK-NEXT: .reg .b64 %rd<2>;1949; CHECK-EMPTY:1950; CHECK-NEXT: // %bb.0:1951; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xi32_param_0];1952; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1953; CHECK-NEXT: add.s32 %r5, %r4, 1;1954; CHECK-NEXT: add.s32 %r6, %r3, 1;1955; CHECK-NEXT: add.s32 %r7, %r2, 1;1956; CHECK-NEXT: add.s32 %r8, %r1, 1;1957; CHECK-NEXT: st.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};1958; CHECK-NEXT: ret;1959 %a.load = load <4 x i32>, ptr addrspace(3) %a1960 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>1961 store <4 x i32> %a.add, ptr addrspace(3) %a1962 ret void1963}1964 1965define void @shared_2xi64(ptr addrspace(3) %a) {1966; CHECK-LABEL: shared_2xi64(1967; CHECK: {1968; CHECK-NEXT: .reg .b64 %rd<6>;1969; CHECK-EMPTY:1970; CHECK-NEXT: // %bb.0:1971; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xi64_param_0];1972; CHECK-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];1973; CHECK-NEXT: add.s64 %rd4, %rd3, 1;1974; CHECK-NEXT: add.s64 %rd5, %rd2, 1;1975; CHECK-NEXT: st.shared.v2.b64 [%rd1], {%rd5, %rd4};1976; CHECK-NEXT: ret;1977 %a.load = load <2 x i64>, ptr addrspace(3) %a1978 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>1979 store <2 x i64> %a.add, ptr addrspace(3) %a1980 ret void1981}1982 1983define void @shared_2xfloat(ptr addrspace(3) %a) {1984; CHECK-LABEL: shared_2xfloat(1985; CHECK: {1986; CHECK-NEXT: .reg .b32 %r<5>;1987; CHECK-NEXT: .reg .b64 %rd<2>;1988; CHECK-EMPTY:1989; CHECK-NEXT: // %bb.0:1990; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xfloat_param_0];1991; CHECK-NEXT: ld.shared.v2.b32 {%r1, %r2}, [%rd1];1992; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;1993; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;1994; CHECK-NEXT: st.shared.v2.b32 [%rd1], {%r4, %r3};1995; CHECK-NEXT: ret;1996 %a.load = load <2 x float>, ptr addrspace(3) %a1997 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>1998 store <2 x float> %a.add, ptr addrspace(3) %a1999 ret void2000}2001 2002define void @shared_4xfloat(ptr addrspace(3) %a) {2003; CHECK-LABEL: shared_4xfloat(2004; CHECK: {2005; CHECK-NEXT: .reg .b32 %r<9>;2006; CHECK-NEXT: .reg .b64 %rd<2>;2007; CHECK-EMPTY:2008; CHECK-NEXT: // %bb.0:2009; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xfloat_param_0];2010; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2011; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;2012; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;2013; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;2014; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;2015; CHECK-NEXT: st.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2016; CHECK-NEXT: ret;2017 %a.load = load <4 x float>, ptr addrspace(3) %a2018 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>2019 store <4 x float> %a.add, ptr addrspace(3) %a2020 ret void2021}2022 2023define void @shared_2xdouble(ptr addrspace(3) %a) {2024; CHECK-LABEL: shared_2xdouble(2025; CHECK: {2026; CHECK-NEXT: .reg .b64 %rd<6>;2027; CHECK-EMPTY:2028; CHECK-NEXT: // %bb.0:2029; CHECK-NEXT: ld.param.b64 %rd1, [shared_2xdouble_param_0];2030; CHECK-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];2031; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;2032; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;2033; CHECK-NEXT: st.shared.v2.b64 [%rd1], {%rd5, %rd4};2034; CHECK-NEXT: ret;2035 %a.load = load <2 x double>, ptr addrspace(3) %a2036 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>2037 store <2 x double> %a.add, ptr addrspace(3) %a2038 ret void2039}2040 2041; shared_volatile2042 2043define void @shared_volatile_2xi8(ptr addrspace(3) %a) {2044; CHECK-LABEL: shared_volatile_2xi8(2045; CHECK: {2046; CHECK-NEXT: .reg .b16 %rs<5>;2047; CHECK-NEXT: .reg .b64 %rd<2>;2048; CHECK-EMPTY:2049; CHECK-NEXT: // %bb.0:2050; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xi8_param_0];2051; CHECK-NEXT: ld.volatile.shared.v2.b8 {%rs1, %rs2}, [%rd1];2052; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2053; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2054; CHECK-NEXT: st.volatile.shared.v2.b8 [%rd1], {%rs4, %rs3};2055; CHECK-NEXT: ret;2056 %a.load = load volatile <2 x i8>, ptr addrspace(3) %a2057 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>2058 store volatile <2 x i8> %a.add, ptr addrspace(3) %a2059 ret void2060}2061 2062define void @shared_volatile_4xi8(ptr addrspace(3) %a) {2063; CHECK-LABEL: shared_volatile_4xi8(2064; CHECK: {2065; CHECK-NEXT: .reg .b16 %rs<9>;2066; CHECK-NEXT: .reg .b32 %r<13>;2067; CHECK-NEXT: .reg .b64 %rd<2>;2068; CHECK-EMPTY:2069; CHECK-NEXT: // %bb.0:2070; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xi8_param_0];2071; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2072; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;2073; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;2074; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2075; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;2076; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;2077; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;2078; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2079; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;2080; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;2081; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;2082; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;2083; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2084; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;2085; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;2086; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;2087; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2088; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;2089; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;2090; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;2091; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r12;2092; CHECK-NEXT: ret;2093 %a.load = load volatile <4 x i8>, ptr addrspace(3) %a2094 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>2095 store volatile <4 x i8> %a.add, ptr addrspace(3) %a2096 ret void2097}2098 2099define void @shared_volatile_8xi8(ptr addrspace(3) %a) {2100; CHECK-LABEL: shared_volatile_8xi8(2101; CHECK: {2102; CHECK-NEXT: .reg .b16 %rs<17>;2103; CHECK-NEXT: .reg .b32 %r<25>;2104; CHECK-NEXT: .reg .b64 %rd<2>;2105; CHECK-EMPTY:2106; CHECK-NEXT: // %bb.0:2107; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_8xi8_param_0];2108; CHECK-NEXT: ld.volatile.shared.v2.b32 {%r1, %r2}, [%rd1];2109; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;2110; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;2111; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2112; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;2113; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;2114; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;2115; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2116; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;2117; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;2118; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;2119; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;2120; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2121; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;2122; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;2123; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;2124; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2125; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;2126; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;2127; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;2128; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;2129; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;2130; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2131; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;2132; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;2133; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;2134; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2135; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;2136; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;2137; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;2138; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;2139; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2140; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;2141; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;2142; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;2143; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2144; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;2145; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;2146; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;2147; CHECK-NEXT: st.volatile.shared.v2.b32 [%rd1], {%r24, %r13};2148; CHECK-NEXT: ret;2149 %a.load = load volatile <8 x i8>, ptr addrspace(3) %a2150 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>2151 store volatile <8 x i8> %a.add, ptr addrspace(3) %a2152 ret void2153}2154 2155define void @shared_volatile_16xi8(ptr addrspace(3) %a) {2156; CHECK-LABEL: shared_volatile_16xi8(2157; CHECK: {2158; CHECK-NEXT: .reg .b16 %rs<33>;2159; CHECK-NEXT: .reg .b32 %r<49>;2160; CHECK-NEXT: .reg .b64 %rd<2>;2161; CHECK-EMPTY:2162; CHECK-NEXT: // %bb.0:2163; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_16xi8_param_0];2164; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2165; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;2166; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;2167; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2168; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;2169; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;2170; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;2171; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2172; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;2173; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;2174; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;2175; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;2176; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2177; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;2178; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;2179; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;2180; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2181; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;2182; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;2183; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;2184; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;2185; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;2186; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2187; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;2188; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;2189; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;2190; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2191; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;2192; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;2193; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;2194; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;2195; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2196; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;2197; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;2198; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;2199; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2200; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;2201; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;2202; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;2203; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;2204; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;2205; CHECK-NEXT: add.s16 %rs18, %rs17, 1;2206; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;2207; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;2208; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;2209; CHECK-NEXT: add.s16 %rs20, %rs19, 1;2210; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;2211; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;2212; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;2213; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;2214; CHECK-NEXT: add.s16 %rs22, %rs21, 1;2215; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;2216; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;2217; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;2218; CHECK-NEXT: add.s16 %rs24, %rs23, 1;2219; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;2220; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;2221; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;2222; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;2223; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;2224; CHECK-NEXT: add.s16 %rs26, %rs25, 1;2225; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;2226; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;2227; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;2228; CHECK-NEXT: add.s16 %rs28, %rs27, 1;2229; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;2230; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;2231; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;2232; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;2233; CHECK-NEXT: add.s16 %rs30, %rs29, 1;2234; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;2235; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;2236; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;2237; CHECK-NEXT: add.s16 %rs32, %rs31, 1;2238; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;2239; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;2240; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;2241; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};2242; CHECK-NEXT: ret;2243 %a.load = load volatile <16 x i8>, ptr addrspace(3) %a2244 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>2245 store volatile <16 x i8> %a.add, ptr addrspace(3) %a2246 ret void2247}2248 2249define void @shared_volatile_2xi16(ptr addrspace(3) %a) {2250; CHECK-LABEL: shared_volatile_2xi16(2251; CHECK: {2252; CHECK-NEXT: .reg .b16 %rs<5>;2253; CHECK-NEXT: .reg .b64 %rd<2>;2254; CHECK-EMPTY:2255; CHECK-NEXT: // %bb.0:2256; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xi16_param_0];2257; CHECK-NEXT: ld.volatile.shared.v2.b16 {%rs1, %rs2}, [%rd1];2258; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2259; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2260; CHECK-NEXT: st.volatile.shared.v2.b16 [%rd1], {%rs4, %rs3};2261; CHECK-NEXT: ret;2262 %a.load = load volatile <2 x i16>, ptr addrspace(3) %a2263 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>2264 store volatile <2 x i16> %a.add, ptr addrspace(3) %a2265 ret void2266}2267 2268define void @shared_volatile_4xi16(ptr addrspace(3) %a) {2269; CHECK-LABEL: shared_volatile_4xi16(2270; CHECK: {2271; CHECK-NEXT: .reg .b16 %rs<9>;2272; CHECK-NEXT: .reg .b64 %rd<2>;2273; CHECK-EMPTY:2274; CHECK-NEXT: // %bb.0:2275; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xi16_param_0];2276; CHECK-NEXT: ld.volatile.shared.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];2277; CHECK-NEXT: add.s16 %rs5, %rs4, 1;2278; CHECK-NEXT: add.s16 %rs6, %rs3, 1;2279; CHECK-NEXT: add.s16 %rs7, %rs2, 1;2280; CHECK-NEXT: add.s16 %rs8, %rs1, 1;2281; CHECK-NEXT: st.volatile.shared.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};2282; CHECK-NEXT: ret;2283 %a.load = load volatile <4 x i16>, ptr addrspace(3) %a2284 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>2285 store volatile <4 x i16> %a.add, ptr addrspace(3) %a2286 ret void2287}2288 2289define void @shared_volatile_8xi16(ptr addrspace(3) %a) {2290; CHECK-LABEL: shared_volatile_8xi16(2291; CHECK: {2292; CHECK-NEXT: .reg .b16 %rs<17>;2293; CHECK-NEXT: .reg .b32 %r<9>;2294; CHECK-NEXT: .reg .b64 %rd<2>;2295; CHECK-EMPTY:2296; CHECK-NEXT: // %bb.0:2297; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_8xi16_param_0];2298; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2299; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;2300; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2301; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2302; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};2303; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;2304; CHECK-NEXT: add.s16 %rs7, %rs6, 1;2305; CHECK-NEXT: add.s16 %rs8, %rs5, 1;2306; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};2307; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;2308; CHECK-NEXT: add.s16 %rs11, %rs10, 1;2309; CHECK-NEXT: add.s16 %rs12, %rs9, 1;2310; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};2311; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;2312; CHECK-NEXT: add.s16 %rs15, %rs14, 1;2313; CHECK-NEXT: add.s16 %rs16, %rs13, 1;2314; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};2315; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2316; CHECK-NEXT: ret;2317 %a.load = load volatile <8 x i16>, ptr addrspace(3) %a2318 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>2319 store volatile <8 x i16> %a.add, ptr addrspace(3) %a2320 ret void2321}2322 2323define void @shared_volatile_2xi32(ptr addrspace(3) %a) {2324; CHECK-LABEL: shared_volatile_2xi32(2325; CHECK: {2326; CHECK-NEXT: .reg .b32 %r<5>;2327; CHECK-NEXT: .reg .b64 %rd<2>;2328; CHECK-EMPTY:2329; CHECK-NEXT: // %bb.0:2330; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xi32_param_0];2331; CHECK-NEXT: ld.volatile.shared.v2.b32 {%r1, %r2}, [%rd1];2332; CHECK-NEXT: add.s32 %r3, %r2, 1;2333; CHECK-NEXT: add.s32 %r4, %r1, 1;2334; CHECK-NEXT: st.volatile.shared.v2.b32 [%rd1], {%r4, %r3};2335; CHECK-NEXT: ret;2336 %a.load = load volatile <2 x i32>, ptr addrspace(3) %a2337 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>2338 store volatile <2 x i32> %a.add, ptr addrspace(3) %a2339 ret void2340}2341 2342define void @shared_volatile_4xi32(ptr addrspace(3) %a) {2343; CHECK-LABEL: shared_volatile_4xi32(2344; CHECK: {2345; CHECK-NEXT: .reg .b32 %r<9>;2346; CHECK-NEXT: .reg .b64 %rd<2>;2347; CHECK-EMPTY:2348; CHECK-NEXT: // %bb.0:2349; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xi32_param_0];2350; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2351; CHECK-NEXT: add.s32 %r5, %r4, 1;2352; CHECK-NEXT: add.s32 %r6, %r3, 1;2353; CHECK-NEXT: add.s32 %r7, %r2, 1;2354; CHECK-NEXT: add.s32 %r8, %r1, 1;2355; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2356; CHECK-NEXT: ret;2357 %a.load = load volatile <4 x i32>, ptr addrspace(3) %a2358 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>2359 store volatile <4 x i32> %a.add, ptr addrspace(3) %a2360 ret void2361}2362 2363define void @shared_volatile_2xi64(ptr addrspace(3) %a) {2364; CHECK-LABEL: shared_volatile_2xi64(2365; CHECK: {2366; CHECK-NEXT: .reg .b64 %rd<6>;2367; CHECK-EMPTY:2368; CHECK-NEXT: // %bb.0:2369; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xi64_param_0];2370; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];2371; CHECK-NEXT: add.s64 %rd4, %rd3, 1;2372; CHECK-NEXT: add.s64 %rd5, %rd2, 1;2373; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd1], {%rd5, %rd4};2374; CHECK-NEXT: ret;2375 %a.load = load volatile <2 x i64>, ptr addrspace(3) %a2376 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>2377 store volatile <2 x i64> %a.add, ptr addrspace(3) %a2378 ret void2379}2380 2381define void @shared_volatile_2xfloat(ptr addrspace(3) %a) {2382; CHECK-LABEL: shared_volatile_2xfloat(2383; CHECK: {2384; CHECK-NEXT: .reg .b32 %r<5>;2385; CHECK-NEXT: .reg .b64 %rd<2>;2386; CHECK-EMPTY:2387; CHECK-NEXT: // %bb.0:2388; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xfloat_param_0];2389; CHECK-NEXT: ld.volatile.shared.v2.b32 {%r1, %r2}, [%rd1];2390; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;2391; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;2392; CHECK-NEXT: st.volatile.shared.v2.b32 [%rd1], {%r4, %r3};2393; CHECK-NEXT: ret;2394 %a.load = load volatile <2 x float>, ptr addrspace(3) %a2395 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>2396 store volatile <2 x float> %a.add, ptr addrspace(3) %a2397 ret void2398}2399 2400define void @shared_volatile_4xfloat(ptr addrspace(3) %a) {2401; CHECK-LABEL: shared_volatile_4xfloat(2402; CHECK: {2403; CHECK-NEXT: .reg .b32 %r<9>;2404; CHECK-NEXT: .reg .b64 %rd<2>;2405; CHECK-EMPTY:2406; CHECK-NEXT: // %bb.0:2407; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xfloat_param_0];2408; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2409; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;2410; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;2411; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;2412; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;2413; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2414; CHECK-NEXT: ret;2415 %a.load = load volatile <4 x float>, ptr addrspace(3) %a2416 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>2417 store volatile <4 x float> %a.add, ptr addrspace(3) %a2418 ret void2419}2420 2421define void @shared_volatile_2xdouble(ptr addrspace(3) %a) {2422; CHECK-LABEL: shared_volatile_2xdouble(2423; CHECK: {2424; CHECK-NEXT: .reg .b64 %rd<6>;2425; CHECK-EMPTY:2426; CHECK-NEXT: // %bb.0:2427; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_2xdouble_param_0];2428; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];2429; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;2430; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;2431; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd1], {%rd5, %rd4};2432; CHECK-NEXT: ret;2433 %a.load = load volatile <2 x double>, ptr addrspace(3) %a2434 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>2435 store volatile <2 x double> %a.add, ptr addrspace(3) %a2436 ret void2437}2438 2439;; local statespace2440 2441; local2442 2443define void @local_2xi8(ptr addrspace(5) %a) {2444; CHECK-LABEL: local_2xi8(2445; CHECK: {2446; CHECK-NEXT: .reg .b16 %rs<5>;2447; CHECK-NEXT: .reg .b64 %rd<2>;2448; CHECK-EMPTY:2449; CHECK-NEXT: // %bb.0:2450; CHECK-NEXT: ld.param.b64 %rd1, [local_2xi8_param_0];2451; CHECK-NEXT: ld.local.v2.b8 {%rs1, %rs2}, [%rd1];2452; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2453; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2454; CHECK-NEXT: st.local.v2.b8 [%rd1], {%rs4, %rs3};2455; CHECK-NEXT: ret;2456 %a.load = load <2 x i8>, ptr addrspace(5) %a2457 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>2458 store <2 x i8> %a.add, ptr addrspace(5) %a2459 ret void2460}2461 2462define void @local_4xi8(ptr addrspace(5) %a) {2463; CHECK-LABEL: local_4xi8(2464; CHECK: {2465; CHECK-NEXT: .reg .b16 %rs<9>;2466; CHECK-NEXT: .reg .b32 %r<13>;2467; CHECK-NEXT: .reg .b64 %rd<2>;2468; CHECK-EMPTY:2469; CHECK-NEXT: // %bb.0:2470; CHECK-NEXT: ld.param.b64 %rd1, [local_4xi8_param_0];2471; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2472; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;2473; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;2474; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2475; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;2476; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;2477; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;2478; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2479; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;2480; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;2481; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;2482; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;2483; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2484; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;2485; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;2486; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;2487; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2488; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;2489; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;2490; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;2491; CHECK-NEXT: st.local.b32 [%rd1], %r12;2492; CHECK-NEXT: ret;2493 %a.load = load <4 x i8>, ptr addrspace(5) %a2494 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>2495 store <4 x i8> %a.add, ptr addrspace(5) %a2496 ret void2497}2498 2499define void @local_8xi8(ptr addrspace(5) %a) {2500; CHECK-LABEL: local_8xi8(2501; CHECK: {2502; CHECK-NEXT: .reg .b16 %rs<17>;2503; CHECK-NEXT: .reg .b32 %r<25>;2504; CHECK-NEXT: .reg .b64 %rd<2>;2505; CHECK-EMPTY:2506; CHECK-NEXT: // %bb.0:2507; CHECK-NEXT: ld.param.b64 %rd1, [local_8xi8_param_0];2508; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];2509; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;2510; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;2511; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2512; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;2513; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;2514; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;2515; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2516; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;2517; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;2518; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;2519; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;2520; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2521; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;2522; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;2523; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;2524; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2525; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;2526; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;2527; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;2528; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;2529; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;2530; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2531; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;2532; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;2533; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;2534; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2535; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;2536; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;2537; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;2538; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;2539; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2540; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;2541; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;2542; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;2543; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2544; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;2545; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;2546; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;2547; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r24, %r13};2548; CHECK-NEXT: ret;2549 %a.load = load <8 x i8>, ptr addrspace(5) %a2550 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>2551 store <8 x i8> %a.add, ptr addrspace(5) %a2552 ret void2553}2554 2555define void @local_16xi8(ptr addrspace(5) %a) {2556; CHECK-LABEL: local_16xi8(2557; CHECK: {2558; CHECK-NEXT: .reg .b16 %rs<33>;2559; CHECK-NEXT: .reg .b32 %r<49>;2560; CHECK-NEXT: .reg .b64 %rd<2>;2561; CHECK-EMPTY:2562; CHECK-NEXT: // %bb.0:2563; CHECK-NEXT: ld.param.b64 %rd1, [local_16xi8_param_0];2564; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2565; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;2566; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;2567; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2568; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;2569; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;2570; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;2571; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2572; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;2573; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;2574; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;2575; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;2576; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2577; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;2578; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;2579; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;2580; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2581; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;2582; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;2583; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;2584; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;2585; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;2586; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2587; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;2588; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;2589; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;2590; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2591; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;2592; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;2593; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;2594; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;2595; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2596; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;2597; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;2598; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;2599; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2600; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;2601; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;2602; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;2603; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;2604; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;2605; CHECK-NEXT: add.s16 %rs18, %rs17, 1;2606; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;2607; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;2608; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;2609; CHECK-NEXT: add.s16 %rs20, %rs19, 1;2610; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;2611; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;2612; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;2613; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;2614; CHECK-NEXT: add.s16 %rs22, %rs21, 1;2615; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;2616; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;2617; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;2618; CHECK-NEXT: add.s16 %rs24, %rs23, 1;2619; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;2620; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;2621; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;2622; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;2623; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;2624; CHECK-NEXT: add.s16 %rs26, %rs25, 1;2625; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;2626; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;2627; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;2628; CHECK-NEXT: add.s16 %rs28, %rs27, 1;2629; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;2630; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;2631; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;2632; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;2633; CHECK-NEXT: add.s16 %rs30, %rs29, 1;2634; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;2635; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;2636; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;2637; CHECK-NEXT: add.s16 %rs32, %rs31, 1;2638; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;2639; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;2640; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;2641; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};2642; CHECK-NEXT: ret;2643 %a.load = load <16 x i8>, ptr addrspace(5) %a2644 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>2645 store <16 x i8> %a.add, ptr addrspace(5) %a2646 ret void2647}2648 2649define void @local_2xi16(ptr addrspace(5) %a) {2650; CHECK-LABEL: local_2xi16(2651; CHECK: {2652; CHECK-NEXT: .reg .b16 %rs<5>;2653; CHECK-NEXT: .reg .b64 %rd<2>;2654; CHECK-EMPTY:2655; CHECK-NEXT: // %bb.0:2656; CHECK-NEXT: ld.param.b64 %rd1, [local_2xi16_param_0];2657; CHECK-NEXT: ld.local.v2.b16 {%rs1, %rs2}, [%rd1];2658; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2659; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2660; CHECK-NEXT: st.local.v2.b16 [%rd1], {%rs4, %rs3};2661; CHECK-NEXT: ret;2662 %a.load = load <2 x i16>, ptr addrspace(5) %a2663 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>2664 store <2 x i16> %a.add, ptr addrspace(5) %a2665 ret void2666}2667 2668define void @local_4xi16(ptr addrspace(5) %a) {2669; CHECK-LABEL: local_4xi16(2670; CHECK: {2671; CHECK-NEXT: .reg .b16 %rs<9>;2672; CHECK-NEXT: .reg .b64 %rd<2>;2673; CHECK-EMPTY:2674; CHECK-NEXT: // %bb.0:2675; CHECK-NEXT: ld.param.b64 %rd1, [local_4xi16_param_0];2676; CHECK-NEXT: ld.local.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];2677; CHECK-NEXT: add.s16 %rs5, %rs4, 1;2678; CHECK-NEXT: add.s16 %rs6, %rs3, 1;2679; CHECK-NEXT: add.s16 %rs7, %rs2, 1;2680; CHECK-NEXT: add.s16 %rs8, %rs1, 1;2681; CHECK-NEXT: st.local.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};2682; CHECK-NEXT: ret;2683 %a.load = load <4 x i16>, ptr addrspace(5) %a2684 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>2685 store <4 x i16> %a.add, ptr addrspace(5) %a2686 ret void2687}2688 2689define void @local_8xi16(ptr addrspace(5) %a) {2690; CHECK-LABEL: local_8xi16(2691; CHECK: {2692; CHECK-NEXT: .reg .b16 %rs<17>;2693; CHECK-NEXT: .reg .b32 %r<9>;2694; CHECK-NEXT: .reg .b64 %rd<2>;2695; CHECK-EMPTY:2696; CHECK-NEXT: // %bb.0:2697; CHECK-NEXT: ld.param.b64 %rd1, [local_8xi16_param_0];2698; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2699; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;2700; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2701; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2702; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};2703; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;2704; CHECK-NEXT: add.s16 %rs7, %rs6, 1;2705; CHECK-NEXT: add.s16 %rs8, %rs5, 1;2706; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};2707; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;2708; CHECK-NEXT: add.s16 %rs11, %rs10, 1;2709; CHECK-NEXT: add.s16 %rs12, %rs9, 1;2710; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};2711; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;2712; CHECK-NEXT: add.s16 %rs15, %rs14, 1;2713; CHECK-NEXT: add.s16 %rs16, %rs13, 1;2714; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};2715; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2716; CHECK-NEXT: ret;2717 %a.load = load <8 x i16>, ptr addrspace(5) %a2718 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>2719 store <8 x i16> %a.add, ptr addrspace(5) %a2720 ret void2721}2722 2723define void @local_2xi32(ptr addrspace(5) %a) {2724; CHECK-LABEL: local_2xi32(2725; CHECK: {2726; CHECK-NEXT: .reg .b32 %r<5>;2727; CHECK-NEXT: .reg .b64 %rd<2>;2728; CHECK-EMPTY:2729; CHECK-NEXT: // %bb.0:2730; CHECK-NEXT: ld.param.b64 %rd1, [local_2xi32_param_0];2731; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];2732; CHECK-NEXT: add.s32 %r3, %r2, 1;2733; CHECK-NEXT: add.s32 %r4, %r1, 1;2734; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r4, %r3};2735; CHECK-NEXT: ret;2736 %a.load = load <2 x i32>, ptr addrspace(5) %a2737 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>2738 store <2 x i32> %a.add, ptr addrspace(5) %a2739 ret void2740}2741 2742define void @local_4xi32(ptr addrspace(5) %a) {2743; CHECK-LABEL: local_4xi32(2744; CHECK: {2745; CHECK-NEXT: .reg .b32 %r<9>;2746; CHECK-NEXT: .reg .b64 %rd<2>;2747; CHECK-EMPTY:2748; CHECK-NEXT: // %bb.0:2749; CHECK-NEXT: ld.param.b64 %rd1, [local_4xi32_param_0];2750; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2751; CHECK-NEXT: add.s32 %r5, %r4, 1;2752; CHECK-NEXT: add.s32 %r6, %r3, 1;2753; CHECK-NEXT: add.s32 %r7, %r2, 1;2754; CHECK-NEXT: add.s32 %r8, %r1, 1;2755; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2756; CHECK-NEXT: ret;2757 %a.load = load <4 x i32>, ptr addrspace(5) %a2758 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>2759 store <4 x i32> %a.add, ptr addrspace(5) %a2760 ret void2761}2762 2763define void @local_2xi64(ptr addrspace(5) %a) {2764; CHECK-LABEL: local_2xi64(2765; CHECK: {2766; CHECK-NEXT: .reg .b64 %rd<6>;2767; CHECK-EMPTY:2768; CHECK-NEXT: // %bb.0:2769; CHECK-NEXT: ld.param.b64 %rd1, [local_2xi64_param_0];2770; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];2771; CHECK-NEXT: add.s64 %rd4, %rd3, 1;2772; CHECK-NEXT: add.s64 %rd5, %rd2, 1;2773; CHECK-NEXT: st.local.v2.b64 [%rd1], {%rd5, %rd4};2774; CHECK-NEXT: ret;2775 %a.load = load <2 x i64>, ptr addrspace(5) %a2776 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>2777 store <2 x i64> %a.add, ptr addrspace(5) %a2778 ret void2779}2780 2781define void @local_2xfloat(ptr addrspace(5) %a) {2782; CHECK-LABEL: local_2xfloat(2783; CHECK: {2784; CHECK-NEXT: .reg .b32 %r<5>;2785; CHECK-NEXT: .reg .b64 %rd<2>;2786; CHECK-EMPTY:2787; CHECK-NEXT: // %bb.0:2788; CHECK-NEXT: ld.param.b64 %rd1, [local_2xfloat_param_0];2789; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];2790; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;2791; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;2792; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r4, %r3};2793; CHECK-NEXT: ret;2794 %a.load = load <2 x float>, ptr addrspace(5) %a2795 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>2796 store <2 x float> %a.add, ptr addrspace(5) %a2797 ret void2798}2799 2800define void @local_4xfloat(ptr addrspace(5) %a) {2801; CHECK-LABEL: local_4xfloat(2802; CHECK: {2803; CHECK-NEXT: .reg .b32 %r<9>;2804; CHECK-NEXT: .reg .b64 %rd<2>;2805; CHECK-EMPTY:2806; CHECK-NEXT: // %bb.0:2807; CHECK-NEXT: ld.param.b64 %rd1, [local_4xfloat_param_0];2808; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2809; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;2810; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;2811; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;2812; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;2813; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};2814; CHECK-NEXT: ret;2815 %a.load = load <4 x float>, ptr addrspace(5) %a2816 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>2817 store <4 x float> %a.add, ptr addrspace(5) %a2818 ret void2819}2820 2821define void @local_2xdouble(ptr addrspace(5) %a) {2822; CHECK-LABEL: local_2xdouble(2823; CHECK: {2824; CHECK-NEXT: .reg .b64 %rd<6>;2825; CHECK-EMPTY:2826; CHECK-NEXT: // %bb.0:2827; CHECK-NEXT: ld.param.b64 %rd1, [local_2xdouble_param_0];2828; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];2829; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;2830; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;2831; CHECK-NEXT: st.local.v2.b64 [%rd1], {%rd5, %rd4};2832; CHECK-NEXT: ret;2833 %a.load = load <2 x double>, ptr addrspace(5) %a2834 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>2835 store <2 x double> %a.add, ptr addrspace(5) %a2836 ret void2837}2838 2839; local_volatile2840 2841define void @local_volatile_2xi8(ptr addrspace(5) %a) {2842; CHECK-LABEL: local_volatile_2xi8(2843; CHECK: {2844; CHECK-NEXT: .reg .b16 %rs<5>;2845; CHECK-NEXT: .reg .b64 %rd<2>;2846; CHECK-EMPTY:2847; CHECK-NEXT: // %bb.0:2848; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xi8_param_0];2849; CHECK-NEXT: ld.local.v2.b8 {%rs1, %rs2}, [%rd1];2850; CHECK-NEXT: add.s16 %rs3, %rs2, 1;2851; CHECK-NEXT: add.s16 %rs4, %rs1, 1;2852; CHECK-NEXT: st.local.v2.b8 [%rd1], {%rs4, %rs3};2853; CHECK-NEXT: ret;2854 %a.load = load volatile <2 x i8>, ptr addrspace(5) %a2855 %a.add = add <2 x i8> %a.load, <i8 1, i8 1>2856 store volatile <2 x i8> %a.add, ptr addrspace(5) %a2857 ret void2858}2859 2860define void @local_volatile_4xi8(ptr addrspace(5) %a) {2861; CHECK-LABEL: local_volatile_4xi8(2862; CHECK: {2863; CHECK-NEXT: .reg .b16 %rs<9>;2864; CHECK-NEXT: .reg .b32 %r<13>;2865; CHECK-NEXT: .reg .b64 %rd<2>;2866; CHECK-EMPTY:2867; CHECK-NEXT: // %bb.0:2868; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xi8_param_0];2869; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2870; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;2871; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;2872; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2873; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;2874; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;2875; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;2876; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2877; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;2878; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;2879; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;2880; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;2881; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2882; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;2883; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;2884; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;2885; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2886; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;2887; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;2888; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;2889; CHECK-NEXT: st.local.b32 [%rd1], %r12;2890; CHECK-NEXT: ret;2891 %a.load = load volatile <4 x i8>, ptr addrspace(5) %a2892 %a.add = add <4 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1>2893 store volatile <4 x i8> %a.add, ptr addrspace(5) %a2894 ret void2895}2896 2897define void @local_volatile_8xi8(ptr addrspace(5) %a) {2898; CHECK-LABEL: local_volatile_8xi8(2899; CHECK: {2900; CHECK-NEXT: .reg .b16 %rs<17>;2901; CHECK-NEXT: .reg .b32 %r<25>;2902; CHECK-NEXT: .reg .b64 %rd<2>;2903; CHECK-EMPTY:2904; CHECK-NEXT: // %bb.0:2905; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_8xi8_param_0];2906; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];2907; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;2908; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;2909; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2910; CHECK-NEXT: cvt.u32.u16 %r4, %rs2;2911; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7772U;2912; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;2913; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2914; CHECK-NEXT: cvt.u32.u16 %r6, %rs4;2915; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 0x3340U;2916; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7771U;2917; CHECK-NEXT: cvt.u16.u32 %rs5, %r8;2918; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2919; CHECK-NEXT: cvt.u32.u16 %r9, %rs6;2920; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;2921; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;2922; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2923; CHECK-NEXT: cvt.u32.u16 %r11, %rs8;2924; CHECK-NEXT: prmt.b32 %r12, %r11, %r9, 0x3340U;2925; CHECK-NEXT: prmt.b32 %r13, %r12, %r7, 0x5410U;2926; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;2927; CHECK-NEXT: cvt.u16.u32 %rs9, %r14;2928; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2929; CHECK-NEXT: cvt.u32.u16 %r15, %rs10;2930; CHECK-NEXT: prmt.b32 %r16, %r1, 0, 0x7772U;2931; CHECK-NEXT: cvt.u16.u32 %rs11, %r16;2932; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2933; CHECK-NEXT: cvt.u32.u16 %r17, %rs12;2934; CHECK-NEXT: prmt.b32 %r18, %r17, %r15, 0x3340U;2935; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;2936; CHECK-NEXT: cvt.u16.u32 %rs13, %r19;2937; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2938; CHECK-NEXT: cvt.u32.u16 %r20, %rs14;2939; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7770U;2940; CHECK-NEXT: cvt.u16.u32 %rs15, %r21;2941; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2942; CHECK-NEXT: cvt.u32.u16 %r22, %rs16;2943; CHECK-NEXT: prmt.b32 %r23, %r22, %r20, 0x3340U;2944; CHECK-NEXT: prmt.b32 %r24, %r23, %r18, 0x5410U;2945; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r24, %r13};2946; CHECK-NEXT: ret;2947 %a.load = load volatile <8 x i8>, ptr addrspace(5) %a2948 %a.add = add <8 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>2949 store volatile <8 x i8> %a.add, ptr addrspace(5) %a2950 ret void2951}2952 2953define void @local_volatile_16xi8(ptr addrspace(5) %a) {2954; CHECK-LABEL: local_volatile_16xi8(2955; CHECK: {2956; CHECK-NEXT: .reg .b16 %rs<33>;2957; CHECK-NEXT: .reg .b32 %r<49>;2958; CHECK-NEXT: .reg .b64 %rd<2>;2959; CHECK-EMPTY:2960; CHECK-NEXT: // %bb.0:2961; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_16xi8_param_0];2962; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];2963; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;2964; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;2965; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2966; CHECK-NEXT: cvt.u32.u16 %r6, %rs2;2967; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7772U;2968; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;2969; CHECK-NEXT: add.s16 %rs4, %rs3, 1;2970; CHECK-NEXT: cvt.u32.u16 %r8, %rs4;2971; CHECK-NEXT: prmt.b32 %r9, %r8, %r6, 0x3340U;2972; CHECK-NEXT: prmt.b32 %r10, %r4, 0, 0x7771U;2973; CHECK-NEXT: cvt.u16.u32 %rs5, %r10;2974; CHECK-NEXT: add.s16 %rs6, %rs5, 1;2975; CHECK-NEXT: cvt.u32.u16 %r11, %rs6;2976; CHECK-NEXT: prmt.b32 %r12, %r4, 0, 0x7770U;2977; CHECK-NEXT: cvt.u16.u32 %rs7, %r12;2978; CHECK-NEXT: add.s16 %rs8, %rs7, 1;2979; CHECK-NEXT: cvt.u32.u16 %r13, %rs8;2980; CHECK-NEXT: prmt.b32 %r14, %r13, %r11, 0x3340U;2981; CHECK-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;2982; CHECK-NEXT: prmt.b32 %r16, %r3, 0, 0x7773U;2983; CHECK-NEXT: cvt.u16.u32 %rs9, %r16;2984; CHECK-NEXT: add.s16 %rs10, %rs9, 1;2985; CHECK-NEXT: cvt.u32.u16 %r17, %rs10;2986; CHECK-NEXT: prmt.b32 %r18, %r3, 0, 0x7772U;2987; CHECK-NEXT: cvt.u16.u32 %rs11, %r18;2988; CHECK-NEXT: add.s16 %rs12, %rs11, 1;2989; CHECK-NEXT: cvt.u32.u16 %r19, %rs12;2990; CHECK-NEXT: prmt.b32 %r20, %r19, %r17, 0x3340U;2991; CHECK-NEXT: prmt.b32 %r21, %r3, 0, 0x7771U;2992; CHECK-NEXT: cvt.u16.u32 %rs13, %r21;2993; CHECK-NEXT: add.s16 %rs14, %rs13, 1;2994; CHECK-NEXT: cvt.u32.u16 %r22, %rs14;2995; CHECK-NEXT: prmt.b32 %r23, %r3, 0, 0x7770U;2996; CHECK-NEXT: cvt.u16.u32 %rs15, %r23;2997; CHECK-NEXT: add.s16 %rs16, %rs15, 1;2998; CHECK-NEXT: cvt.u32.u16 %r24, %rs16;2999; CHECK-NEXT: prmt.b32 %r25, %r24, %r22, 0x3340U;3000; CHECK-NEXT: prmt.b32 %r26, %r25, %r20, 0x5410U;3001; CHECK-NEXT: prmt.b32 %r27, %r2, 0, 0x7773U;3002; CHECK-NEXT: cvt.u16.u32 %rs17, %r27;3003; CHECK-NEXT: add.s16 %rs18, %rs17, 1;3004; CHECK-NEXT: cvt.u32.u16 %r28, %rs18;3005; CHECK-NEXT: prmt.b32 %r29, %r2, 0, 0x7772U;3006; CHECK-NEXT: cvt.u16.u32 %rs19, %r29;3007; CHECK-NEXT: add.s16 %rs20, %rs19, 1;3008; CHECK-NEXT: cvt.u32.u16 %r30, %rs20;3009; CHECK-NEXT: prmt.b32 %r31, %r30, %r28, 0x3340U;3010; CHECK-NEXT: prmt.b32 %r32, %r2, 0, 0x7771U;3011; CHECK-NEXT: cvt.u16.u32 %rs21, %r32;3012; CHECK-NEXT: add.s16 %rs22, %rs21, 1;3013; CHECK-NEXT: cvt.u32.u16 %r33, %rs22;3014; CHECK-NEXT: prmt.b32 %r34, %r2, 0, 0x7770U;3015; CHECK-NEXT: cvt.u16.u32 %rs23, %r34;3016; CHECK-NEXT: add.s16 %rs24, %rs23, 1;3017; CHECK-NEXT: cvt.u32.u16 %r35, %rs24;3018; CHECK-NEXT: prmt.b32 %r36, %r35, %r33, 0x3340U;3019; CHECK-NEXT: prmt.b32 %r37, %r36, %r31, 0x5410U;3020; CHECK-NEXT: prmt.b32 %r38, %r1, 0, 0x7773U;3021; CHECK-NEXT: cvt.u16.u32 %rs25, %r38;3022; CHECK-NEXT: add.s16 %rs26, %rs25, 1;3023; CHECK-NEXT: cvt.u32.u16 %r39, %rs26;3024; CHECK-NEXT: prmt.b32 %r40, %r1, 0, 0x7772U;3025; CHECK-NEXT: cvt.u16.u32 %rs27, %r40;3026; CHECK-NEXT: add.s16 %rs28, %rs27, 1;3027; CHECK-NEXT: cvt.u32.u16 %r41, %rs28;3028; CHECK-NEXT: prmt.b32 %r42, %r41, %r39, 0x3340U;3029; CHECK-NEXT: prmt.b32 %r43, %r1, 0, 0x7771U;3030; CHECK-NEXT: cvt.u16.u32 %rs29, %r43;3031; CHECK-NEXT: add.s16 %rs30, %rs29, 1;3032; CHECK-NEXT: cvt.u32.u16 %r44, %rs30;3033; CHECK-NEXT: prmt.b32 %r45, %r1, 0, 0x7770U;3034; CHECK-NEXT: cvt.u16.u32 %rs31, %r45;3035; CHECK-NEXT: add.s16 %rs32, %rs31, 1;3036; CHECK-NEXT: cvt.u32.u16 %r46, %rs32;3037; CHECK-NEXT: prmt.b32 %r47, %r46, %r44, 0x3340U;3038; CHECK-NEXT: prmt.b32 %r48, %r47, %r42, 0x5410U;3039; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r48, %r37, %r26, %r15};3040; CHECK-NEXT: ret;3041 %a.load = load volatile <16 x i8>, ptr addrspace(5) %a3042 %a.add = add <16 x i8> %a.load, <i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1, i8 1>3043 store volatile <16 x i8> %a.add, ptr addrspace(5) %a3044 ret void3045}3046 3047define void @local_volatile_2xi16(ptr addrspace(5) %a) {3048; CHECK-LABEL: local_volatile_2xi16(3049; CHECK: {3050; CHECK-NEXT: .reg .b16 %rs<5>;3051; CHECK-NEXT: .reg .b64 %rd<2>;3052; CHECK-EMPTY:3053; CHECK-NEXT: // %bb.0:3054; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xi16_param_0];3055; CHECK-NEXT: ld.local.v2.b16 {%rs1, %rs2}, [%rd1];3056; CHECK-NEXT: add.s16 %rs3, %rs2, 1;3057; CHECK-NEXT: add.s16 %rs4, %rs1, 1;3058; CHECK-NEXT: st.local.v2.b16 [%rd1], {%rs4, %rs3};3059; CHECK-NEXT: ret;3060 %a.load = load volatile <2 x i16>, ptr addrspace(5) %a3061 %a.add = add <2 x i16> %a.load, <i16 1, i16 1>3062 store volatile <2 x i16> %a.add, ptr addrspace(5) %a3063 ret void3064}3065 3066define void @local_volatile_4xi16(ptr addrspace(5) %a) {3067; CHECK-LABEL: local_volatile_4xi16(3068; CHECK: {3069; CHECK-NEXT: .reg .b16 %rs<9>;3070; CHECK-NEXT: .reg .b64 %rd<2>;3071; CHECK-EMPTY:3072; CHECK-NEXT: // %bb.0:3073; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xi16_param_0];3074; CHECK-NEXT: ld.local.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];3075; CHECK-NEXT: add.s16 %rs5, %rs4, 1;3076; CHECK-NEXT: add.s16 %rs6, %rs3, 1;3077; CHECK-NEXT: add.s16 %rs7, %rs2, 1;3078; CHECK-NEXT: add.s16 %rs8, %rs1, 1;3079; CHECK-NEXT: st.local.v4.b16 [%rd1], {%rs8, %rs7, %rs6, %rs5};3080; CHECK-NEXT: ret;3081 %a.load = load volatile <4 x i16>, ptr addrspace(5) %a3082 %a.add = add <4 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1>3083 store volatile <4 x i16> %a.add, ptr addrspace(5) %a3084 ret void3085}3086 3087define void @local_volatile_8xi16(ptr addrspace(5) %a) {3088; CHECK-LABEL: local_volatile_8xi16(3089; CHECK: {3090; CHECK-NEXT: .reg .b16 %rs<17>;3091; CHECK-NEXT: .reg .b32 %r<9>;3092; CHECK-NEXT: .reg .b64 %rd<2>;3093; CHECK-EMPTY:3094; CHECK-NEXT: // %bb.0:3095; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_8xi16_param_0];3096; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];3097; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4;3098; CHECK-NEXT: add.s16 %rs3, %rs2, 1;3099; CHECK-NEXT: add.s16 %rs4, %rs1, 1;3100; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};3101; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r3;3102; CHECK-NEXT: add.s16 %rs7, %rs6, 1;3103; CHECK-NEXT: add.s16 %rs8, %rs5, 1;3104; CHECK-NEXT: mov.b32 %r6, {%rs8, %rs7};3105; CHECK-NEXT: mov.b32 {%rs9, %rs10}, %r2;3106; CHECK-NEXT: add.s16 %rs11, %rs10, 1;3107; CHECK-NEXT: add.s16 %rs12, %rs9, 1;3108; CHECK-NEXT: mov.b32 %r7, {%rs12, %rs11};3109; CHECK-NEXT: mov.b32 {%rs13, %rs14}, %r1;3110; CHECK-NEXT: add.s16 %rs15, %rs14, 1;3111; CHECK-NEXT: add.s16 %rs16, %rs13, 1;3112; CHECK-NEXT: mov.b32 %r8, {%rs16, %rs15};3113; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};3114; CHECK-NEXT: ret;3115 %a.load = load volatile <8 x i16>, ptr addrspace(5) %a3116 %a.add = add <8 x i16> %a.load, <i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1, i16 1>3117 store volatile <8 x i16> %a.add, ptr addrspace(5) %a3118 ret void3119}3120 3121define void @local_volatile_2xi32(ptr addrspace(5) %a) {3122; CHECK-LABEL: local_volatile_2xi32(3123; CHECK: {3124; CHECK-NEXT: .reg .b32 %r<5>;3125; CHECK-NEXT: .reg .b64 %rd<2>;3126; CHECK-EMPTY:3127; CHECK-NEXT: // %bb.0:3128; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xi32_param_0];3129; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];3130; CHECK-NEXT: add.s32 %r3, %r2, 1;3131; CHECK-NEXT: add.s32 %r4, %r1, 1;3132; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r4, %r3};3133; CHECK-NEXT: ret;3134 %a.load = load volatile <2 x i32>, ptr addrspace(5) %a3135 %a.add = add <2 x i32> %a.load, <i32 1, i32 1>3136 store volatile <2 x i32> %a.add, ptr addrspace(5) %a3137 ret void3138}3139 3140define void @local_volatile_4xi32(ptr addrspace(5) %a) {3141; CHECK-LABEL: local_volatile_4xi32(3142; CHECK: {3143; CHECK-NEXT: .reg .b32 %r<9>;3144; CHECK-NEXT: .reg .b64 %rd<2>;3145; CHECK-EMPTY:3146; CHECK-NEXT: // %bb.0:3147; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xi32_param_0];3148; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];3149; CHECK-NEXT: add.s32 %r5, %r4, 1;3150; CHECK-NEXT: add.s32 %r6, %r3, 1;3151; CHECK-NEXT: add.s32 %r7, %r2, 1;3152; CHECK-NEXT: add.s32 %r8, %r1, 1;3153; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};3154; CHECK-NEXT: ret;3155 %a.load = load volatile <4 x i32>, ptr addrspace(5) %a3156 %a.add = add <4 x i32> %a.load, <i32 1, i32 1, i32 1, i32 1>3157 store volatile <4 x i32> %a.add, ptr addrspace(5) %a3158 ret void3159}3160 3161define void @local_volatile_2xi64(ptr addrspace(5) %a) {3162; CHECK-LABEL: local_volatile_2xi64(3163; CHECK: {3164; CHECK-NEXT: .reg .b64 %rd<6>;3165; CHECK-EMPTY:3166; CHECK-NEXT: // %bb.0:3167; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xi64_param_0];3168; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];3169; CHECK-NEXT: add.s64 %rd4, %rd3, 1;3170; CHECK-NEXT: add.s64 %rd5, %rd2, 1;3171; CHECK-NEXT: st.local.v2.b64 [%rd1], {%rd5, %rd4};3172; CHECK-NEXT: ret;3173 %a.load = load volatile <2 x i64>, ptr addrspace(5) %a3174 %a.add = add <2 x i64> %a.load, <i64 1, i64 1>3175 store volatile <2 x i64> %a.add, ptr addrspace(5) %a3176 ret void3177}3178 3179define void @local_volatile_2xfloat(ptr addrspace(5) %a) {3180; CHECK-LABEL: local_volatile_2xfloat(3181; CHECK: {3182; CHECK-NEXT: .reg .b32 %r<5>;3183; CHECK-NEXT: .reg .b64 %rd<2>;3184; CHECK-EMPTY:3185; CHECK-NEXT: // %bb.0:3186; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xfloat_param_0];3187; CHECK-NEXT: ld.local.v2.b32 {%r1, %r2}, [%rd1];3188; CHECK-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;3189; CHECK-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;3190; CHECK-NEXT: st.local.v2.b32 [%rd1], {%r4, %r3};3191; CHECK-NEXT: ret;3192 %a.load = load volatile <2 x float>, ptr addrspace(5) %a3193 %a.add = fadd <2 x float> %a.load, <float 1., float 1.>3194 store volatile <2 x float> %a.add, ptr addrspace(5) %a3195 ret void3196}3197 3198define void @local_volatile_4xfloat(ptr addrspace(5) %a) {3199; CHECK-LABEL: local_volatile_4xfloat(3200; CHECK: {3201; CHECK-NEXT: .reg .b32 %r<9>;3202; CHECK-NEXT: .reg .b64 %rd<2>;3203; CHECK-EMPTY:3204; CHECK-NEXT: // %bb.0:3205; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xfloat_param_0];3206; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];3207; CHECK-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;3208; CHECK-NEXT: add.rn.f32 %r6, %r3, 0f3F800000;3209; CHECK-NEXT: add.rn.f32 %r7, %r2, 0f3F800000;3210; CHECK-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;3211; CHECK-NEXT: st.local.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};3212; CHECK-NEXT: ret;3213 %a.load = load volatile <4 x float>, ptr addrspace(5) %a3214 %a.add = fadd <4 x float> %a.load, <float 1., float 1., float 1., float 1.>3215 store volatile <4 x float> %a.add, ptr addrspace(5) %a3216 ret void3217}3218 3219define void @local_volatile_2xdouble(ptr addrspace(5) %a) {3220; CHECK-LABEL: local_volatile_2xdouble(3221; CHECK: {3222; CHECK-NEXT: .reg .b64 %rd<6>;3223; CHECK-EMPTY:3224; CHECK-NEXT: // %bb.0:3225; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_2xdouble_param_0];3226; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];3227; CHECK-NEXT: add.rn.f64 %rd4, %rd3, 0d3FF0000000000000;3228; CHECK-NEXT: add.rn.f64 %rd5, %rd2, 0d3FF0000000000000;3229; CHECK-NEXT: st.local.v2.b64 [%rd1], {%rd5, %rd4};3230; CHECK-NEXT: ret;3231 %a.load = load volatile <2 x double>, ptr addrspace(5) %a3232 %a.add = fadd <2 x double> %a.load, <double 1., double 1.>3233 store volatile <2 x double> %a.add, ptr addrspace(5) %a3234 ret void3235}3236