614 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s | FileCheck -check-prefix=ENABLED %s3; RUN: llc -disable-nvptx-load-store-vectorizer < %s | FileCheck -check-prefix=DISABLED %s4; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}5; RUN: %if ptxas %{ llc -disable-nvptx-load-store-vectorizer < %s | %ptxas-verify %}6 7target triple = "nvptx64-nvidia-cuda"8 9; Check that the load-store vectorizer is enabled by default for nvptx, and10; that it's disabled by the appropriate flag.11 12define i32 @f(ptr %p) {13; ENABLED-LABEL: f(14; ENABLED: {15; ENABLED-NEXT: .reg .b32 %r<4>;16; ENABLED-NEXT: .reg .b64 %rd<2>;17; ENABLED-EMPTY:18; ENABLED-NEXT: // %bb.0:19; ENABLED-NEXT: ld.param.b64 %rd1, [f_param_0];20; ENABLED-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];21; ENABLED-NEXT: add.s32 %r3, %r1, %r2;22; ENABLED-NEXT: st.param.b32 [func_retval0], %r3;23; ENABLED-NEXT: ret;24;25; DISABLED-LABEL: f(26; DISABLED: {27; DISABLED-NEXT: .reg .b32 %r<4>;28; DISABLED-NEXT: .reg .b64 %rd<2>;29; DISABLED-EMPTY:30; DISABLED-NEXT: // %bb.0:31; DISABLED-NEXT: ld.param.b64 %rd1, [f_param_0];32; DISABLED-NEXT: ld.b32 %r1, [%rd1];33; DISABLED-NEXT: ld.b32 %r2, [%rd1+4];34; DISABLED-NEXT: add.s32 %r3, %r1, %r2;35; DISABLED-NEXT: st.param.b32 [func_retval0], %r3;36; DISABLED-NEXT: ret;37 %p.1 = getelementptr i32, ptr %p, i32 138 %v0 = load i32, ptr %p, align 839 %v1 = load i32, ptr %p.1, align 440 %sum = add i32 %v0, %v141 ret i32 %sum42}43 44define half @fh(ptr %p) {45; ENABLED-LABEL: fh(46; ENABLED: {47; ENABLED-NEXT: .reg .b16 %rs<10>;48; ENABLED-NEXT: .reg .b32 %r<13>;49; ENABLED-NEXT: .reg .b64 %rd<2>;50; ENABLED-EMPTY:51; ENABLED-NEXT: // %bb.0:52; ENABLED-NEXT: ld.param.b64 %rd1, [fh_param_0];53; ENABLED-NEXT: ld.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];54; ENABLED-NEXT: ld.b16 %rs5, [%rd1+8];55; ENABLED-NEXT: cvt.f32.f16 %r1, %rs2;56; ENABLED-NEXT: cvt.f32.f16 %r2, %rs1;57; ENABLED-NEXT: add.rn.f32 %r3, %r2, %r1;58; ENABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;59; ENABLED-NEXT: cvt.f32.f16 %r4, %rs4;60; ENABLED-NEXT: cvt.f32.f16 %r5, %rs3;61; ENABLED-NEXT: add.rn.f32 %r6, %r5, %r4;62; ENABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;63; ENABLED-NEXT: cvt.f32.f16 %r7, %rs7;64; ENABLED-NEXT: cvt.f32.f16 %r8, %rs6;65; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r7;66; ENABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;67; ENABLED-NEXT: cvt.f32.f16 %r10, %rs8;68; ENABLED-NEXT: cvt.f32.f16 %r11, %rs5;69; ENABLED-NEXT: add.rn.f32 %r12, %r10, %r11;70; ENABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;71; ENABLED-NEXT: st.param.b16 [func_retval0], %rs9;72; ENABLED-NEXT: ret;73;74; DISABLED-LABEL: fh(75; DISABLED: {76; DISABLED-NEXT: .reg .b16 %rs<10>;77; DISABLED-NEXT: .reg .b32 %r<13>;78; DISABLED-NEXT: .reg .b64 %rd<2>;79; DISABLED-EMPTY:80; DISABLED-NEXT: // %bb.0:81; DISABLED-NEXT: ld.param.b64 %rd1, [fh_param_0];82; DISABLED-NEXT: ld.b16 %rs1, [%rd1];83; DISABLED-NEXT: ld.b16 %rs2, [%rd1+2];84; DISABLED-NEXT: ld.b16 %rs3, [%rd1+4];85; DISABLED-NEXT: ld.b16 %rs4, [%rd1+6];86; DISABLED-NEXT: ld.b16 %rs5, [%rd1+8];87; DISABLED-NEXT: cvt.f32.f16 %r1, %rs2;88; DISABLED-NEXT: cvt.f32.f16 %r2, %rs1;89; DISABLED-NEXT: add.rn.f32 %r3, %r2, %r1;90; DISABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;91; DISABLED-NEXT: cvt.f32.f16 %r4, %rs4;92; DISABLED-NEXT: cvt.f32.f16 %r5, %rs3;93; DISABLED-NEXT: add.rn.f32 %r6, %r5, %r4;94; DISABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;95; DISABLED-NEXT: cvt.f32.f16 %r7, %rs7;96; DISABLED-NEXT: cvt.f32.f16 %r8, %rs6;97; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r7;98; DISABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;99; DISABLED-NEXT: cvt.f32.f16 %r10, %rs8;100; DISABLED-NEXT: cvt.f32.f16 %r11, %rs5;101; DISABLED-NEXT: add.rn.f32 %r12, %r10, %r11;102; DISABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;103; DISABLED-NEXT: st.param.b16 [func_retval0], %rs9;104; DISABLED-NEXT: ret;105 %p.1 = getelementptr half, ptr %p, i32 1106 %p.2 = getelementptr half, ptr %p, i32 2107 %p.3 = getelementptr half, ptr %p, i32 3108 %p.4 = getelementptr half, ptr %p, i32 4109 %v0 = load half, ptr %p, align 64110 %v1 = load half, ptr %p.1, align 4111 %v2 = load half, ptr %p.2, align 4112 %v3 = load half, ptr %p.3, align 4113 %v4 = load half, ptr %p.4, align 4114 %sum1 = fadd half %v0, %v1115 %sum2 = fadd half %v2, %v3116 %sum3 = fadd half %sum1, %sum2117 %sum = fadd half %sum3, %v4118 ret half %sum119}120 121define float @ff(ptr %p) {122; ENABLED-LABEL: ff(123; ENABLED: {124; ENABLED-NEXT: .reg .b32 %r<10>;125; ENABLED-NEXT: .reg .b64 %rd<2>;126; ENABLED-EMPTY:127; ENABLED-NEXT: // %bb.0:128; ENABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];129; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];130; ENABLED-NEXT: ld.b32 %r5, [%rd1+16];131; ENABLED-NEXT: add.rn.f32 %r6, %r1, %r2;132; ENABLED-NEXT: add.rn.f32 %r7, %r3, %r4;133; ENABLED-NEXT: add.rn.f32 %r8, %r6, %r7;134; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r5;135; ENABLED-NEXT: st.param.b32 [func_retval0], %r9;136; ENABLED-NEXT: ret;137;138; DISABLED-LABEL: ff(139; DISABLED: {140; DISABLED-NEXT: .reg .b32 %r<10>;141; DISABLED-NEXT: .reg .b64 %rd<2>;142; DISABLED-EMPTY:143; DISABLED-NEXT: // %bb.0:144; DISABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];145; DISABLED-NEXT: ld.b32 %r1, [%rd1];146; DISABLED-NEXT: ld.b32 %r2, [%rd1+4];147; DISABLED-NEXT: ld.b32 %r3, [%rd1+8];148; DISABLED-NEXT: ld.b32 %r4, [%rd1+12];149; DISABLED-NEXT: ld.b32 %r5, [%rd1+16];150; DISABLED-NEXT: add.rn.f32 %r6, %r1, %r2;151; DISABLED-NEXT: add.rn.f32 %r7, %r3, %r4;152; DISABLED-NEXT: add.rn.f32 %r8, %r6, %r7;153; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r5;154; DISABLED-NEXT: st.param.b32 [func_retval0], %r9;155; DISABLED-NEXT: ret;156 %p.1 = getelementptr float, ptr %p, i32 1157 %p.2 = getelementptr float, ptr %p, i32 2158 %p.3 = getelementptr float, ptr %p, i32 3159 %p.4 = getelementptr float, ptr %p, i32 4160 %v0 = load float, ptr %p, align 64161 %v1 = load float, ptr %p.1, align 4162 %v2 = load float, ptr %p.2, align 4163 %v3 = load float, ptr %p.3, align 4164 %v4 = load float, ptr %p.4, align 4165 %sum1 = fadd float %v0, %v1166 %sum2 = fadd float %v2, %v3167 %sum3 = fadd float %sum1, %sum2168 %sum = fadd float %sum3, %v4169 ret float %sum170}171 172define void @combine_v16i8(ptr noundef align 16 %ptr1, ptr noundef align 16 %ptr2) {173; ENABLED-LABEL: combine_v16i8(174; ENABLED: {175; ENABLED-NEXT: .reg .b32 %r<36>;176; ENABLED-NEXT: .reg .b64 %rd<3>;177; ENABLED-EMPTY:178; ENABLED-NEXT: // %bb.0:179; ENABLED-NEXT: ld.param.b64 %rd1, [combine_v16i8_param_0];180; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];181; ENABLED-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;182; ENABLED-NEXT: prmt.b32 %r6, %r4, 0, 0x7772U;183; ENABLED-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;184; ENABLED-NEXT: prmt.b32 %r8, %r3, 0, 0x7773U;185; ENABLED-NEXT: prmt.b32 %r9, %r3, 0, 0x7772U;186; ENABLED-NEXT: prmt.b32 %r10, %r3, 0, 0x7771U;187; ENABLED-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U;188; ENABLED-NEXT: prmt.b32 %r12, %r2, 0, 0x7772U;189; ENABLED-NEXT: prmt.b32 %r13, %r2, 0, 0x7771U;190; ENABLED-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;191; ENABLED-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;192; ENABLED-NEXT: prmt.b32 %r16, %r1, 0, 0x7771U;193; ENABLED-NEXT: ld.param.b64 %rd2, [combine_v16i8_param_1];194; ENABLED-NEXT: and.b32 %r17, %r1, 255;195; ENABLED-NEXT: and.b32 %r18, %r2, 255;196; ENABLED-NEXT: and.b32 %r19, %r3, 255;197; ENABLED-NEXT: and.b32 %r20, %r4, 255;198; ENABLED-NEXT: add.s32 %r21, %r17, %r16;199; ENABLED-NEXT: add.s32 %r22, %r21, %r15;200; ENABLED-NEXT: add.s32 %r23, %r22, %r14;201; ENABLED-NEXT: add.s32 %r24, %r23, %r18;202; ENABLED-NEXT: add.s32 %r25, %r24, %r13;203; ENABLED-NEXT: add.s32 %r26, %r25, %r12;204; ENABLED-NEXT: add.s32 %r27, %r26, %r11;205; ENABLED-NEXT: add.s32 %r28, %r27, %r19;206; ENABLED-NEXT: add.s32 %r29, %r28, %r10;207; ENABLED-NEXT: add.s32 %r30, %r29, %r9;208; ENABLED-NEXT: add.s32 %r31, %r30, %r8;209; ENABLED-NEXT: add.s32 %r32, %r31, %r20;210; ENABLED-NEXT: add.s32 %r33, %r32, %r7;211; ENABLED-NEXT: add.s32 %r34, %r33, %r6;212; ENABLED-NEXT: add.s32 %r35, %r34, %r5;213; ENABLED-NEXT: st.b32 [%rd2], %r35;214; ENABLED-NEXT: ret;215;216; DISABLED-LABEL: combine_v16i8(217; DISABLED: {218; DISABLED-NEXT: .reg .b32 %r<32>;219; DISABLED-NEXT: .reg .b64 %rd<3>;220; DISABLED-EMPTY:221; DISABLED-NEXT: // %bb.0:222; DISABLED-NEXT: ld.param.b64 %rd1, [combine_v16i8_param_0];223; DISABLED-NEXT: ld.b8 %r1, [%rd1];224; DISABLED-NEXT: ld.param.b64 %rd2, [combine_v16i8_param_1];225; DISABLED-NEXT: ld.b8 %r2, [%rd1+1];226; DISABLED-NEXT: ld.b8 %r3, [%rd1+2];227; DISABLED-NEXT: ld.b8 %r4, [%rd1+3];228; DISABLED-NEXT: ld.b8 %r5, [%rd1+4];229; DISABLED-NEXT: ld.b8 %r6, [%rd1+5];230; DISABLED-NEXT: ld.b8 %r7, [%rd1+6];231; DISABLED-NEXT: ld.b8 %r8, [%rd1+7];232; DISABLED-NEXT: ld.b8 %r9, [%rd1+8];233; DISABLED-NEXT: ld.b8 %r10, [%rd1+9];234; DISABLED-NEXT: ld.b8 %r11, [%rd1+10];235; DISABLED-NEXT: ld.b8 %r12, [%rd1+11];236; DISABLED-NEXT: ld.b8 %r13, [%rd1+12];237; DISABLED-NEXT: ld.b8 %r14, [%rd1+13];238; DISABLED-NEXT: ld.b8 %r15, [%rd1+14];239; DISABLED-NEXT: ld.b8 %r16, [%rd1+15];240; DISABLED-NEXT: add.s32 %r17, %r1, %r2;241; DISABLED-NEXT: add.s32 %r18, %r17, %r3;242; DISABLED-NEXT: add.s32 %r19, %r18, %r4;243; DISABLED-NEXT: add.s32 %r20, %r19, %r5;244; DISABLED-NEXT: add.s32 %r21, %r20, %r6;245; DISABLED-NEXT: add.s32 %r22, %r21, %r7;246; DISABLED-NEXT: add.s32 %r23, %r22, %r8;247; DISABLED-NEXT: add.s32 %r24, %r23, %r9;248; DISABLED-NEXT: add.s32 %r25, %r24, %r10;249; DISABLED-NEXT: add.s32 %r26, %r25, %r11;250; DISABLED-NEXT: add.s32 %r27, %r26, %r12;251; DISABLED-NEXT: add.s32 %r28, %r27, %r13;252; DISABLED-NEXT: add.s32 %r29, %r28, %r14;253; DISABLED-NEXT: add.s32 %r30, %r29, %r15;254; DISABLED-NEXT: add.s32 %r31, %r30, %r16;255; DISABLED-NEXT: st.b32 [%rd2], %r31;256; DISABLED-NEXT: ret;257 %val0 = load i8, ptr %ptr1, align 16258 %ptr1.1 = getelementptr inbounds i8, ptr %ptr1, i64 1259 %val1 = load i8, ptr %ptr1.1, align 1260 %ptr1.2 = getelementptr inbounds i8, ptr %ptr1, i64 2261 %val2 = load i8, ptr %ptr1.2, align 2262 %ptr1.3 = getelementptr inbounds i8, ptr %ptr1, i64 3263 %val3 = load i8, ptr %ptr1.3, align 1264 %ptr1.4 = getelementptr inbounds i8, ptr %ptr1, i64 4265 %val4 = load i8, ptr %ptr1.4, align 4266 %ptr1.5 = getelementptr inbounds i8, ptr %ptr1, i64 5267 %val5 = load i8, ptr %ptr1.5, align 1268 %ptr1.6 = getelementptr inbounds i8, ptr %ptr1, i64 6269 %val6 = load i8, ptr %ptr1.6, align 2270 %ptr1.7 = getelementptr inbounds i8, ptr %ptr1, i64 7271 %val7 = load i8, ptr %ptr1.7, align 1272 %ptr1.8 = getelementptr inbounds i8, ptr %ptr1, i64 8273 %val8 = load i8, ptr %ptr1.8, align 8274 %ptr1.9 = getelementptr inbounds i8, ptr %ptr1, i64 9275 %val9 = load i8, ptr %ptr1.9, align 1276 %ptr1.10 = getelementptr inbounds i8, ptr %ptr1, i64 10277 %val10 = load i8, ptr %ptr1.10, align 2278 %ptr1.11 = getelementptr inbounds i8, ptr %ptr1, i64 11279 %val11 = load i8, ptr %ptr1.11, align 1280 %ptr1.12 = getelementptr inbounds i8, ptr %ptr1, i64 12281 %val12 = load i8, ptr %ptr1.12, align 4282 %ptr1.13 = getelementptr inbounds i8, ptr %ptr1, i64 13283 %val13 = load i8, ptr %ptr1.13, align 1284 %ptr1.14 = getelementptr inbounds i8, ptr %ptr1, i64 14285 %val14 = load i8, ptr %ptr1.14, align 2286 %ptr1.15 = getelementptr inbounds i8, ptr %ptr1, i64 15287 %val15 = load i8, ptr %ptr1.15, align 1288 %lane0 = zext i8 %val0 to i32289 %lane1 = zext i8 %val1 to i32290 %lane2 = zext i8 %val2 to i32291 %lane3 = zext i8 %val3 to i32292 %lane4 = zext i8 %val4 to i32293 %lane5 = zext i8 %val5 to i32294 %lane6 = zext i8 %val6 to i32295 %lane7 = zext i8 %val7 to i32296 %lane8 = zext i8 %val8 to i32297 %lane9 = zext i8 %val9 to i32298 %lane10 = zext i8 %val10 to i32299 %lane11 = zext i8 %val11 to i32300 %lane12 = zext i8 %val12 to i32301 %lane13 = zext i8 %val13 to i32302 %lane14 = zext i8 %val14 to i32303 %lane15 = zext i8 %val15 to i32304 %red.1 = add i32 %lane0, %lane1305 %red.2 = add i32 %red.1, %lane2306 %red.3 = add i32 %red.2, %lane3307 %red.4 = add i32 %red.3, %lane4308 %red.5 = add i32 %red.4, %lane5309 %red.6 = add i32 %red.5, %lane6310 %red.7 = add i32 %red.6, %lane7311 %red.8 = add i32 %red.7, %lane8312 %red.9 = add i32 %red.8, %lane9313 %red.10 = add i32 %red.9, %lane10314 %red.11 = add i32 %red.10, %lane11315 %red.12 = add i32 %red.11, %lane12316 %red.13 = add i32 %red.12, %lane13317 %red.14 = add i32 %red.13, %lane14318 %red = add i32 %red.14, %lane15319 store i32 %red, ptr %ptr2, align 4320 ret void321}322 323define void @combine_v16i8_unaligned(ptr noundef align 8 %ptr1, ptr noundef align 16 %ptr2) {324; ENABLED-LABEL: combine_v16i8_unaligned(325; ENABLED: {326; ENABLED-NEXT: .reg .b32 %r<36>;327; ENABLED-NEXT: .reg .b64 %rd<3>;328; ENABLED-EMPTY:329; ENABLED-NEXT: // %bb.0:330; ENABLED-NEXT: ld.param.b64 %rd1, [combine_v16i8_unaligned_param_0];331; ENABLED-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];332; ENABLED-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;333; ENABLED-NEXT: prmt.b32 %r4, %r2, 0, 0x7772U;334; ENABLED-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;335; ENABLED-NEXT: prmt.b32 %r6, %r1, 0, 0x7773U;336; ENABLED-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;337; ENABLED-NEXT: prmt.b32 %r8, %r1, 0, 0x7771U;338; ENABLED-NEXT: ld.param.b64 %rd2, [combine_v16i8_unaligned_param_1];339; ENABLED-NEXT: ld.v2.b32 {%r9, %r10}, [%rd1+8];340; ENABLED-NEXT: prmt.b32 %r11, %r10, 0, 0x7773U;341; ENABLED-NEXT: prmt.b32 %r12, %r10, 0, 0x7772U;342; ENABLED-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U;343; ENABLED-NEXT: prmt.b32 %r14, %r9, 0, 0x7773U;344; ENABLED-NEXT: prmt.b32 %r15, %r9, 0, 0x7772U;345; ENABLED-NEXT: prmt.b32 %r16, %r9, 0, 0x7771U;346; ENABLED-NEXT: and.b32 %r17, %r1, 255;347; ENABLED-NEXT: and.b32 %r18, %r2, 255;348; ENABLED-NEXT: and.b32 %r19, %r9, 255;349; ENABLED-NEXT: and.b32 %r20, %r10, 255;350; ENABLED-NEXT: add.s32 %r21, %r17, %r8;351; ENABLED-NEXT: add.s32 %r22, %r21, %r7;352; ENABLED-NEXT: add.s32 %r23, %r22, %r6;353; ENABLED-NEXT: add.s32 %r24, %r23, %r18;354; ENABLED-NEXT: add.s32 %r25, %r24, %r5;355; ENABLED-NEXT: add.s32 %r26, %r25, %r4;356; ENABLED-NEXT: add.s32 %r27, %r26, %r3;357; ENABLED-NEXT: add.s32 %r28, %r27, %r19;358; ENABLED-NEXT: add.s32 %r29, %r28, %r16;359; ENABLED-NEXT: add.s32 %r30, %r29, %r15;360; ENABLED-NEXT: add.s32 %r31, %r30, %r14;361; ENABLED-NEXT: add.s32 %r32, %r31, %r20;362; ENABLED-NEXT: add.s32 %r33, %r32, %r13;363; ENABLED-NEXT: add.s32 %r34, %r33, %r12;364; ENABLED-NEXT: add.s32 %r35, %r34, %r11;365; ENABLED-NEXT: st.b32 [%rd2], %r35;366; ENABLED-NEXT: ret;367;368; DISABLED-LABEL: combine_v16i8_unaligned(369; DISABLED: {370; DISABLED-NEXT: .reg .b32 %r<32>;371; DISABLED-NEXT: .reg .b64 %rd<3>;372; DISABLED-EMPTY:373; DISABLED-NEXT: // %bb.0:374; DISABLED-NEXT: ld.param.b64 %rd1, [combine_v16i8_unaligned_param_0];375; DISABLED-NEXT: ld.b8 %r1, [%rd1];376; DISABLED-NEXT: ld.param.b64 %rd2, [combine_v16i8_unaligned_param_1];377; DISABLED-NEXT: ld.b8 %r2, [%rd1+1];378; DISABLED-NEXT: ld.b8 %r3, [%rd1+2];379; DISABLED-NEXT: ld.b8 %r4, [%rd1+3];380; DISABLED-NEXT: ld.b8 %r5, [%rd1+4];381; DISABLED-NEXT: ld.b8 %r6, [%rd1+5];382; DISABLED-NEXT: ld.b8 %r7, [%rd1+6];383; DISABLED-NEXT: ld.b8 %r8, [%rd1+7];384; DISABLED-NEXT: ld.b8 %r9, [%rd1+8];385; DISABLED-NEXT: ld.b8 %r10, [%rd1+9];386; DISABLED-NEXT: ld.b8 %r11, [%rd1+10];387; DISABLED-NEXT: ld.b8 %r12, [%rd1+11];388; DISABLED-NEXT: ld.b8 %r13, [%rd1+12];389; DISABLED-NEXT: ld.b8 %r14, [%rd1+13];390; DISABLED-NEXT: ld.b8 %r15, [%rd1+14];391; DISABLED-NEXT: ld.b8 %r16, [%rd1+15];392; DISABLED-NEXT: add.s32 %r17, %r1, %r2;393; DISABLED-NEXT: add.s32 %r18, %r17, %r3;394; DISABLED-NEXT: add.s32 %r19, %r18, %r4;395; DISABLED-NEXT: add.s32 %r20, %r19, %r5;396; DISABLED-NEXT: add.s32 %r21, %r20, %r6;397; DISABLED-NEXT: add.s32 %r22, %r21, %r7;398; DISABLED-NEXT: add.s32 %r23, %r22, %r8;399; DISABLED-NEXT: add.s32 %r24, %r23, %r9;400; DISABLED-NEXT: add.s32 %r25, %r24, %r10;401; DISABLED-NEXT: add.s32 %r26, %r25, %r11;402; DISABLED-NEXT: add.s32 %r27, %r26, %r12;403; DISABLED-NEXT: add.s32 %r28, %r27, %r13;404; DISABLED-NEXT: add.s32 %r29, %r28, %r14;405; DISABLED-NEXT: add.s32 %r30, %r29, %r15;406; DISABLED-NEXT: add.s32 %r31, %r30, %r16;407; DISABLED-NEXT: st.b32 [%rd2], %r31;408; DISABLED-NEXT: ret;409 %val0 = load i8, ptr %ptr1, align 8410 %ptr1.1 = getelementptr inbounds i8, ptr %ptr1, i64 1411 %val1 = load i8, ptr %ptr1.1, align 1412 %ptr1.2 = getelementptr inbounds i8, ptr %ptr1, i64 2413 %val2 = load i8, ptr %ptr1.2, align 2414 %ptr1.3 = getelementptr inbounds i8, ptr %ptr1, i64 3415 %val3 = load i8, ptr %ptr1.3, align 1416 %ptr1.4 = getelementptr inbounds i8, ptr %ptr1, i64 4417 %val4 = load i8, ptr %ptr1.4, align 4418 %ptr1.5 = getelementptr inbounds i8, ptr %ptr1, i64 5419 %val5 = load i8, ptr %ptr1.5, align 1420 %ptr1.6 = getelementptr inbounds i8, ptr %ptr1, i64 6421 %val6 = load i8, ptr %ptr1.6, align 2422 %ptr1.7 = getelementptr inbounds i8, ptr %ptr1, i64 7423 %val7 = load i8, ptr %ptr1.7, align 1424 %ptr1.8 = getelementptr inbounds i8, ptr %ptr1, i64 8425 %val8 = load i8, ptr %ptr1.8, align 8426 %ptr1.9 = getelementptr inbounds i8, ptr %ptr1, i64 9427 %val9 = load i8, ptr %ptr1.9, align 1428 %ptr1.10 = getelementptr inbounds i8, ptr %ptr1, i64 10429 %val10 = load i8, ptr %ptr1.10, align 2430 %ptr1.11 = getelementptr inbounds i8, ptr %ptr1, i64 11431 %val11 = load i8, ptr %ptr1.11, align 1432 %ptr1.12 = getelementptr inbounds i8, ptr %ptr1, i64 12433 %val12 = load i8, ptr %ptr1.12, align 4434 %ptr1.13 = getelementptr inbounds i8, ptr %ptr1, i64 13435 %val13 = load i8, ptr %ptr1.13, align 1436 %ptr1.14 = getelementptr inbounds i8, ptr %ptr1, i64 14437 %val14 = load i8, ptr %ptr1.14, align 2438 %ptr1.15 = getelementptr inbounds i8, ptr %ptr1, i64 15439 %val15 = load i8, ptr %ptr1.15, align 1440 %lane0 = zext i8 %val0 to i32441 %lane1 = zext i8 %val1 to i32442 %lane2 = zext i8 %val2 to i32443 %lane3 = zext i8 %val3 to i32444 %lane4 = zext i8 %val4 to i32445 %lane5 = zext i8 %val5 to i32446 %lane6 = zext i8 %val6 to i32447 %lane7 = zext i8 %val7 to i32448 %lane8 = zext i8 %val8 to i32449 %lane9 = zext i8 %val9 to i32450 %lane10 = zext i8 %val10 to i32451 %lane11 = zext i8 %val11 to i32452 %lane12 = zext i8 %val12 to i32453 %lane13 = zext i8 %val13 to i32454 %lane14 = zext i8 %val14 to i32455 %lane15 = zext i8 %val15 to i32456 %red.1 = add i32 %lane0, %lane1457 %red.2 = add i32 %red.1, %lane2458 %red.3 = add i32 %red.2, %lane3459 %red.4 = add i32 %red.3, %lane4460 %red.5 = add i32 %red.4, %lane5461 %red.6 = add i32 %red.5, %lane6462 %red.7 = add i32 %red.6, %lane7463 %red.8 = add i32 %red.7, %lane8464 %red.9 = add i32 %red.8, %lane9465 %red.10 = add i32 %red.9, %lane10466 %red.11 = add i32 %red.10, %lane11467 %red.12 = add i32 %red.11, %lane12468 %red.13 = add i32 %red.12, %lane13469 %red.14 = add i32 %red.13, %lane14470 %red = add i32 %red.14, %lane15471 store i32 %red, ptr %ptr2, align 4472 ret void473}474 475 476define void @combine_v8i16(ptr noundef align 16 %ptr1, ptr noundef align 16 %ptr2) {477; ENABLED-LABEL: combine_v8i16(478; ENABLED: {479; ENABLED-NEXT: .reg .b16 %rs<9>;480; ENABLED-NEXT: .reg .b32 %r<20>;481; ENABLED-NEXT: .reg .b64 %rd<3>;482; ENABLED-EMPTY:483; ENABLED-NEXT: // %bb.0:484; ENABLED-NEXT: ld.param.b64 %rd1, [combine_v8i16_param_0];485; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];486; ENABLED-NEXT: mov.b32 {%rs1, %rs2}, %r4;487; ENABLED-NEXT: mov.b32 {%rs3, %rs4}, %r3;488; ENABLED-NEXT: mov.b32 {%rs5, %rs6}, %r2;489; ENABLED-NEXT: mov.b32 {%rs7, %rs8}, %r1;490; ENABLED-NEXT: ld.param.b64 %rd2, [combine_v8i16_param_1];491; ENABLED-NEXT: cvt.u32.u16 %r5, %rs7;492; ENABLED-NEXT: cvt.u32.u16 %r6, %rs8;493; ENABLED-NEXT: cvt.u32.u16 %r7, %rs5;494; ENABLED-NEXT: cvt.u32.u16 %r8, %rs6;495; ENABLED-NEXT: cvt.u32.u16 %r9, %rs3;496; ENABLED-NEXT: cvt.u32.u16 %r10, %rs4;497; ENABLED-NEXT: cvt.u32.u16 %r11, %rs1;498; ENABLED-NEXT: cvt.u32.u16 %r12, %rs2;499; ENABLED-NEXT: add.s32 %r13, %r5, %r6;500; ENABLED-NEXT: add.s32 %r14, %r13, %r7;501; ENABLED-NEXT: add.s32 %r15, %r14, %r8;502; ENABLED-NEXT: add.s32 %r16, %r15, %r9;503; ENABLED-NEXT: add.s32 %r17, %r16, %r10;504; ENABLED-NEXT: add.s32 %r18, %r17, %r11;505; ENABLED-NEXT: add.s32 %r19, %r18, %r12;506; ENABLED-NEXT: st.b32 [%rd2], %r19;507; ENABLED-NEXT: ret;508;509; DISABLED-LABEL: combine_v8i16(510; DISABLED: {511; DISABLED-NEXT: .reg .b32 %r<16>;512; DISABLED-NEXT: .reg .b64 %rd<3>;513; DISABLED-EMPTY:514; DISABLED-NEXT: // %bb.0:515; DISABLED-NEXT: ld.param.b64 %rd1, [combine_v8i16_param_0];516; DISABLED-NEXT: ld.b16 %r1, [%rd1];517; DISABLED-NEXT: ld.param.b64 %rd2, [combine_v8i16_param_1];518; DISABLED-NEXT: ld.b16 %r2, [%rd1+2];519; DISABLED-NEXT: ld.b16 %r3, [%rd1+4];520; DISABLED-NEXT: ld.b16 %r4, [%rd1+6];521; DISABLED-NEXT: ld.b16 %r5, [%rd1+8];522; DISABLED-NEXT: ld.b16 %r6, [%rd1+10];523; DISABLED-NEXT: ld.b16 %r7, [%rd1+12];524; DISABLED-NEXT: ld.b16 %r8, [%rd1+14];525; DISABLED-NEXT: add.s32 %r9, %r1, %r2;526; DISABLED-NEXT: add.s32 %r10, %r9, %r3;527; DISABLED-NEXT: add.s32 %r11, %r10, %r4;528; DISABLED-NEXT: add.s32 %r12, %r11, %r5;529; DISABLED-NEXT: add.s32 %r13, %r12, %r6;530; DISABLED-NEXT: add.s32 %r14, %r13, %r7;531; DISABLED-NEXT: add.s32 %r15, %r14, %r8;532; DISABLED-NEXT: st.b32 [%rd2], %r15;533; DISABLED-NEXT: ret;534 %val0 = load i16, ptr %ptr1, align 16535 %ptr1.1 = getelementptr inbounds i16, ptr %ptr1, i64 1536 %val1 = load i16, ptr %ptr1.1, align 2537 %ptr1.2 = getelementptr inbounds i16, ptr %ptr1, i64 2538 %val2 = load i16, ptr %ptr1.2, align 4539 %ptr1.3 = getelementptr inbounds i16, ptr %ptr1, i64 3540 %val3 = load i16, ptr %ptr1.3, align 2541 %ptr1.4 = getelementptr inbounds i16, ptr %ptr1, i64 4542 %val4 = load i16, ptr %ptr1.4, align 4543 %ptr1.5 = getelementptr inbounds i16, ptr %ptr1, i64 5544 %val5 = load i16, ptr %ptr1.5, align 2545 %ptr1.6 = getelementptr inbounds i16, ptr %ptr1, i64 6546 %val6 = load i16, ptr %ptr1.6, align 4547 %ptr1.7 = getelementptr inbounds i16, ptr %ptr1, i64 7548 %val7 = load i16, ptr %ptr1.7, align 2549 %lane0 = zext i16 %val0 to i32550 %lane1 = zext i16 %val1 to i32551 %lane2 = zext i16 %val2 to i32552 %lane3 = zext i16 %val3 to i32553 %lane4 = zext i16 %val4 to i32554 %lane5 = zext i16 %val5 to i32555 %lane6 = zext i16 %val6 to i32556 %lane7 = zext i16 %val7 to i32557 %red.1 = add i32 %lane0, %lane1558 %red.2 = add i32 %red.1, %lane2559 %red.3 = add i32 %red.2, %lane3560 %red.4 = add i32 %red.3, %lane4561 %red.5 = add i32 %red.4, %lane5562 %red.6 = add i32 %red.5, %lane6563 %red = add i32 %red.6, %lane7564 store i32 %red, ptr %ptr2, align 4565 ret void566}567 568define void @combine_v4i32(ptr noundef align 16 %ptr1, ptr noundef align 16 %ptr2) {569; ENABLED-LABEL: combine_v4i32(570; ENABLED: {571; ENABLED-NEXT: .reg .b32 %r<8>;572; ENABLED-NEXT: .reg .b64 %rd<3>;573; ENABLED-EMPTY:574; ENABLED-NEXT: // %bb.0:575; ENABLED-NEXT: ld.param.b64 %rd1, [combine_v4i32_param_0];576; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];577; ENABLED-NEXT: ld.param.b64 %rd2, [combine_v4i32_param_1];578; ENABLED-NEXT: add.s32 %r5, %r1, %r2;579; ENABLED-NEXT: add.s32 %r6, %r5, %r3;580; ENABLED-NEXT: add.s32 %r7, %r6, %r4;581; ENABLED-NEXT: st.b32 [%rd2], %r7;582; ENABLED-NEXT: ret;583;584; DISABLED-LABEL: combine_v4i32(585; DISABLED: {586; DISABLED-NEXT: .reg .b32 %r<8>;587; DISABLED-NEXT: .reg .b64 %rd<3>;588; DISABLED-EMPTY:589; DISABLED-NEXT: // %bb.0:590; DISABLED-NEXT: ld.param.b64 %rd1, [combine_v4i32_param_0];591; DISABLED-NEXT: ld.b32 %r1, [%rd1];592; DISABLED-NEXT: ld.param.b64 %rd2, [combine_v4i32_param_1];593; DISABLED-NEXT: ld.b32 %r2, [%rd1+4];594; DISABLED-NEXT: ld.b32 %r3, [%rd1+8];595; DISABLED-NEXT: ld.b32 %r4, [%rd1+12];596; DISABLED-NEXT: add.s32 %r5, %r1, %r2;597; DISABLED-NEXT: add.s32 %r6, %r5, %r3;598; DISABLED-NEXT: add.s32 %r7, %r6, %r4;599; DISABLED-NEXT: st.b32 [%rd2], %r7;600; DISABLED-NEXT: ret;601 %val0 = load i32, ptr %ptr1, align 16602 %ptr1.1 = getelementptr inbounds i32, ptr %ptr1, i64 1603 %val1 = load i32, ptr %ptr1.1, align 4604 %ptr1.2 = getelementptr inbounds i32, ptr %ptr1, i64 2605 %val2 = load i32, ptr %ptr1.2, align 8606 %ptr1.3 = getelementptr inbounds i32, ptr %ptr1, i64 3607 %val3 = load i32, ptr %ptr1.3, align 4608 %red.1 = add i32 %val0, %val1609 %red.2 = add i32 %red.1, %val2610 %red = add i32 %red.2, %val3611 store i32 %red, ptr %ptr2, align 4612 ret void613}614