739 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 62; RUN: llc < %s -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 | %ptxas-verify %}4target triple = "nvptx64-nvidia-cuda"5 6; Test dynamic insertelt at the beginning of a chain7define <4 x i32> @dynamic_at_beginning(i32 %idx) {8; CHECK-LABEL: dynamic_at_beginning(9; CHECK: {10; CHECK-NEXT: .local .align 4 .b8 __local_depot0[16];11; CHECK-NEXT: .reg .b64 %SP;12; CHECK-NEXT: .reg .b64 %SPL;13; CHECK-NEXT: .reg .b32 %r<3>;14; CHECK-NEXT: .reg .b64 %rd<6>;15; CHECK-EMPTY:16; CHECK-NEXT: // %bb.0:17; CHECK-NEXT: mov.b64 %SPL, __local_depot0;18; CHECK-NEXT: cvta.local.u64 %SP, %SPL;19; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_at_beginning_param_0];20; CHECK-NEXT: and.b64 %rd2, %rd1, 3;21; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;22; CHECK-NEXT: add.u64 %rd4, %SP, 0;23; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;24; CHECK-NEXT: st.b32 [%rd5], 10;25; CHECK-NEXT: ld.b32 %r1, [%SP+12];26; CHECK-NEXT: ld.b32 %r2, [%SP];27; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r2, 20, 30, %r1};28; CHECK-NEXT: ret;29 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx30 %v1 = insertelement <4 x i32> %v0, i32 20, i32 131 %v2 = insertelement <4 x i32> %v1, i32 30, i32 232 ret <4 x i32> %v233}34 35; Test dynamic insertelt at the end of a chain36define <4 x i32> @dynamic_at_end(i32 %idx) {37; CHECK-LABEL: dynamic_at_end(38; CHECK: {39; CHECK-NEXT: .local .align 4 .b8 __local_depot1[16];40; CHECK-NEXT: .reg .b64 %SP;41; CHECK-NEXT: .reg .b64 %SPL;42; CHECK-NEXT: .reg .b32 %r<5>;43; CHECK-NEXT: .reg .b64 %rd<6>;44; CHECK-EMPTY:45; CHECK-NEXT: // %bb.0:46; CHECK-NEXT: mov.b64 %SPL, __local_depot1;47; CHECK-NEXT: cvta.local.u64 %SP, %SPL;48; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_at_end_param_0];49; CHECK-NEXT: and.b64 %rd2, %rd1, 3;50; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;51; CHECK-NEXT: add.u64 %rd4, %SP, 0;52; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;53; CHECK-NEXT: st.b32 [%SP+4], 20;54; CHECK-NEXT: st.b32 [%SP], 10;55; CHECK-NEXT: st.b32 [%rd5], 30;56; CHECK-NEXT: ld.b32 %r1, [%SP+12];57; CHECK-NEXT: ld.b32 %r2, [%SP+8];58; CHECK-NEXT: ld.b32 %r3, [%SP+4];59; CHECK-NEXT: ld.b32 %r4, [%SP];60; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};61; CHECK-NEXT: ret;62 %v0 = insertelement <4 x i32> poison, i32 10, i32 063 %v1 = insertelement <4 x i32> %v0, i32 20, i32 164 %v2 = insertelement <4 x i32> %v1, i32 30, i32 %idx65 ret <4 x i32> %v266}67 68; Test dynamic insertelt in the middle of a chain69define <4 x i32> @dynamic_in_middle(i32 %idx) {70; CHECK-LABEL: dynamic_in_middle(71; CHECK: {72; CHECK-NEXT: .local .align 4 .b8 __local_depot2[16];73; CHECK-NEXT: .reg .b64 %SP;74; CHECK-NEXT: .reg .b64 %SPL;75; CHECK-NEXT: .reg .b32 %r<4>;76; CHECK-NEXT: .reg .b64 %rd<6>;77; CHECK-EMPTY:78; CHECK-NEXT: // %bb.0:79; CHECK-NEXT: mov.b64 %SPL, __local_depot2;80; CHECK-NEXT: cvta.local.u64 %SP, %SPL;81; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_in_middle_param_0];82; CHECK-NEXT: and.b64 %rd2, %rd1, 3;83; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;84; CHECK-NEXT: add.u64 %rd4, %SP, 0;85; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;86; CHECK-NEXT: st.b32 [%SP], 10;87; CHECK-NEXT: st.b32 [%rd5], 20;88; CHECK-NEXT: ld.b32 %r1, [%SP+12];89; CHECK-NEXT: ld.b32 %r2, [%SP+4];90; CHECK-NEXT: ld.b32 %r3, [%SP];91; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r3, %r2, 30, %r1};92; CHECK-NEXT: ret;93 %v0 = insertelement <4 x i32> poison, i32 10, i32 094 %v1 = insertelement <4 x i32> %v0, i32 20, i32 %idx95 %v2 = insertelement <4 x i32> %v1, i32 30, i32 296 ret <4 x i32> %v297}98 99; Test repeated dynamic insertelt with the same index100define <4 x i32> @repeated_same_index(i32 %idx) {101; CHECK-LABEL: repeated_same_index(102; CHECK: {103; CHECK-NEXT: .local .align 4 .b8 __local_depot3[16];104; CHECK-NEXT: .reg .b64 %SP;105; CHECK-NEXT: .reg .b64 %SPL;106; CHECK-NEXT: .reg .b32 %r<5>;107; CHECK-NEXT: .reg .b64 %rd<6>;108; CHECK-EMPTY:109; CHECK-NEXT: // %bb.0:110; CHECK-NEXT: mov.b64 %SPL, __local_depot3;111; CHECK-NEXT: cvta.local.u64 %SP, %SPL;112; CHECK-NEXT: ld.param.b32 %rd1, [repeated_same_index_param_0];113; CHECK-NEXT: and.b64 %rd2, %rd1, 3;114; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;115; CHECK-NEXT: add.u64 %rd4, %SP, 0;116; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;117; CHECK-NEXT: st.b32 [%rd5], 20;118; CHECK-NEXT: ld.b32 %r1, [%SP+12];119; CHECK-NEXT: ld.b32 %r2, [%SP+8];120; CHECK-NEXT: ld.b32 %r3, [%SP+4];121; CHECK-NEXT: ld.b32 %r4, [%SP];122; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};123; CHECK-NEXT: ret;124 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx125 %v1 = insertelement <4 x i32> %v0, i32 20, i32 %idx126 ret <4 x i32> %v1127}128 129; Test multiple dynamic insertelts130define <4 x i32> @multiple_dynamic(i32 %idx0, i32 %idx1) {131; CHECK-LABEL: multiple_dynamic(132; CHECK: {133; CHECK-NEXT: .local .align 4 .b8 __local_depot4[16];134; CHECK-NEXT: .reg .b64 %SP;135; CHECK-NEXT: .reg .b64 %SPL;136; CHECK-NEXT: .reg .b32 %r<5>;137; CHECK-NEXT: .reg .b64 %rd<10>;138; CHECK-EMPTY:139; CHECK-NEXT: // %bb.0:140; CHECK-NEXT: mov.b64 %SPL, __local_depot4;141; CHECK-NEXT: cvta.local.u64 %SP, %SPL;142; CHECK-NEXT: ld.param.b32 %rd1, [multiple_dynamic_param_0];143; CHECK-NEXT: and.b64 %rd2, %rd1, 3;144; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;145; CHECK-NEXT: add.u64 %rd4, %SP, 0;146; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;147; CHECK-NEXT: st.b32 [%rd5], 10;148; CHECK-NEXT: ld.param.b32 %rd6, [multiple_dynamic_param_1];149; CHECK-NEXT: and.b64 %rd7, %rd6, 3;150; CHECK-NEXT: shl.b64 %rd8, %rd7, 2;151; CHECK-NEXT: add.s64 %rd9, %rd4, %rd8;152; CHECK-NEXT: st.b32 [%rd9], 20;153; CHECK-NEXT: ld.b32 %r1, [%SP+12];154; CHECK-NEXT: ld.b32 %r2, [%SP+8];155; CHECK-NEXT: ld.b32 %r3, [%SP+4];156; CHECK-NEXT: ld.b32 %r4, [%SP];157; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};158; CHECK-NEXT: ret;159 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx0160 %v1 = insertelement <4 x i32> %v0, i32 20, i32 %idx1161 ret <4 x i32> %v1162}163 164; Test chain with all dynamic insertelts165define <4 x i32> @all_dynamic(i32 %idx0, i32 %idx1, i32 %idx2, i32 %idx3) {166; CHECK-LABEL: all_dynamic(167; CHECK: {168; CHECK-NEXT: .local .align 4 .b8 __local_depot5[16];169; CHECK-NEXT: .reg .b64 %SP;170; CHECK-NEXT: .reg .b64 %SPL;171; CHECK-NEXT: .reg .b32 %r<5>;172; CHECK-NEXT: .reg .b64 %rd<18>;173; CHECK-EMPTY:174; CHECK-NEXT: // %bb.0:175; CHECK-NEXT: mov.b64 %SPL, __local_depot5;176; CHECK-NEXT: cvta.local.u64 %SP, %SPL;177; CHECK-NEXT: ld.param.b32 %rd1, [all_dynamic_param_0];178; CHECK-NEXT: and.b64 %rd2, %rd1, 3;179; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;180; CHECK-NEXT: add.u64 %rd4, %SP, 0;181; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;182; CHECK-NEXT: ld.param.b32 %rd6, [all_dynamic_param_1];183; CHECK-NEXT: and.b64 %rd7, %rd6, 3;184; CHECK-NEXT: shl.b64 %rd8, %rd7, 2;185; CHECK-NEXT: add.s64 %rd9, %rd4, %rd8;186; CHECK-NEXT: ld.param.b32 %rd10, [all_dynamic_param_2];187; CHECK-NEXT: and.b64 %rd11, %rd10, 3;188; CHECK-NEXT: shl.b64 %rd12, %rd11, 2;189; CHECK-NEXT: add.s64 %rd13, %rd4, %rd12;190; CHECK-NEXT: st.b32 [%rd5], 10;191; CHECK-NEXT: st.b32 [%rd9], 20;192; CHECK-NEXT: st.b32 [%rd13], 30;193; CHECK-NEXT: ld.param.b32 %rd14, [all_dynamic_param_3];194; CHECK-NEXT: and.b64 %rd15, %rd14, 3;195; CHECK-NEXT: shl.b64 %rd16, %rd15, 2;196; CHECK-NEXT: add.s64 %rd17, %rd4, %rd16;197; CHECK-NEXT: st.b32 [%rd17], 40;198; CHECK-NEXT: ld.b32 %r1, [%SP+12];199; CHECK-NEXT: ld.b32 %r2, [%SP+8];200; CHECK-NEXT: ld.b32 %r3, [%SP+4];201; CHECK-NEXT: ld.b32 %r4, [%SP];202; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};203; CHECK-NEXT: ret;204 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx0205 %v1 = insertelement <4 x i32> %v0, i32 20, i32 %idx1206 %v2 = insertelement <4 x i32> %v1, i32 30, i32 %idx2207 %v3 = insertelement <4 x i32> %v2, i32 40, i32 %idx3208 ret <4 x i32> %v3209}210 211; Test mixed constant and dynamic insertelts with high ratio of dynamic ones.212; Should lower all insertelts to stores.213define <4 x i32> @mix_dynamic_constant(i32 %idx0, i32 %idx1) {214; CHECK-LABEL: mix_dynamic_constant(215; CHECK: {216; CHECK-NEXT: .local .align 4 .b8 __local_depot6[16];217; CHECK-NEXT: .reg .b64 %SP;218; CHECK-NEXT: .reg .b64 %SPL;219; CHECK-NEXT: .reg .b32 %r<5>;220; CHECK-NEXT: .reg .b64 %rd<10>;221; CHECK-EMPTY:222; CHECK-NEXT: // %bb.0:223; CHECK-NEXT: mov.b64 %SPL, __local_depot6;224; CHECK-NEXT: cvta.local.u64 %SP, %SPL;225; CHECK-NEXT: ld.param.b32 %rd1, [mix_dynamic_constant_param_0];226; CHECK-NEXT: and.b64 %rd2, %rd1, 3;227; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;228; CHECK-NEXT: add.u64 %rd4, %SP, 0;229; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;230; CHECK-NEXT: st.b32 [%rd5], 10;231; CHECK-NEXT: ld.param.b32 %rd6, [mix_dynamic_constant_param_1];232; CHECK-NEXT: and.b64 %rd7, %rd6, 3;233; CHECK-NEXT: shl.b64 %rd8, %rd7, 2;234; CHECK-NEXT: add.s64 %rd9, %rd4, %rd8;235; CHECK-NEXT: st.b32 [%SP+4], 20;236; CHECK-NEXT: st.b32 [%rd9], 30;237; CHECK-NEXT: ld.b32 %r1, [%SP+12];238; CHECK-NEXT: ld.b32 %r2, [%SP+8];239; CHECK-NEXT: ld.b32 %r3, [%SP+4];240; CHECK-NEXT: ld.b32 %r4, [%SP];241; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};242; CHECK-NEXT: ret;243 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx0244 %v1 = insertelement <4 x i32> %v0, i32 20, i32 1245 %v2 = insertelement <4 x i32> %v1, i32 30, i32 %idx1246 ret <4 x i32> %v2247}248 249; Test two separate chains that don't interfere250define void @two_separate_chains(i32 %idx0, i32 %idx1, ptr %out0, ptr %out1) {251; CHECK-LABEL: two_separate_chains(252; CHECK: {253; CHECK-NEXT: .local .align 4 .b8 __local_depot7[32];254; CHECK-NEXT: .reg .b64 %SP;255; CHECK-NEXT: .reg .b64 %SPL;256; CHECK-NEXT: .reg .b32 %r<7>;257; CHECK-NEXT: .reg .b64 %rd<13>;258; CHECK-EMPTY:259; CHECK-NEXT: // %bb.0:260; CHECK-NEXT: mov.b64 %SPL, __local_depot7;261; CHECK-NEXT: cvta.local.u64 %SP, %SPL;262; CHECK-NEXT: ld.param.b32 %rd1, [two_separate_chains_param_0];263; CHECK-NEXT: and.b64 %rd2, %rd1, 3;264; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;265; CHECK-NEXT: add.u64 %rd4, %SP, 16;266; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;267; CHECK-NEXT: st.b32 [%rd5], 10;268; CHECK-NEXT: ld.param.b32 %rd6, [two_separate_chains_param_1];269; CHECK-NEXT: and.b64 %rd7, %rd6, 3;270; CHECK-NEXT: shl.b64 %rd8, %rd7, 2;271; CHECK-NEXT: add.u64 %rd9, %SP, 0;272; CHECK-NEXT: add.s64 %rd10, %rd9, %rd8;273; CHECK-NEXT: ld.b32 %r1, [%SP+28];274; CHECK-NEXT: ld.b32 %r2, [%SP+24];275; CHECK-NEXT: ld.b32 %r3, [%SP+16];276; CHECK-NEXT: ld.param.b64 %rd11, [two_separate_chains_param_2];277; CHECK-NEXT: st.b32 [%rd10], 30;278; CHECK-NEXT: ld.param.b64 %rd12, [two_separate_chains_param_3];279; CHECK-NEXT: ld.b32 %r4, [%SP+12];280; CHECK-NEXT: ld.b32 %r5, [%SP+4];281; CHECK-NEXT: ld.b32 %r6, [%SP];282; CHECK-NEXT: st.v4.b32 [%rd11], {%r3, 20, %r2, %r1};283; CHECK-NEXT: st.v4.b32 [%rd12], {%r6, %r5, 40, %r4};284; CHECK-NEXT: ret;285 ; Chain 1286 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx0287 %v1 = insertelement <4 x i32> %v0, i32 20, i32 1288 289 ; Chain 2290 %w0 = insertelement <4 x i32> poison, i32 30, i32 %idx1291 %w1 = insertelement <4 x i32> %w0, i32 40, i32 2292 293 store <4 x i32> %v1, ptr %out0294 store <4 x i32> %w1, ptr %out1295 ret void296}297 298; Test overlapping chains (chain 2 starts from middle of chain 1)299define void @overlapping_chains(i32 %idx0, i32 %idx1, ptr %out0, ptr %out1) {300; CHECK-LABEL: overlapping_chains(301; CHECK: {302; CHECK-NEXT: .local .align 4 .b8 __local_depot8[32];303; CHECK-NEXT: .reg .b64 %SP;304; CHECK-NEXT: .reg .b64 %SPL;305; CHECK-NEXT: .reg .b32 %r<7>;306; CHECK-NEXT: .reg .b64 %rd<14>;307; CHECK-EMPTY:308; CHECK-NEXT: // %bb.0:309; CHECK-NEXT: mov.b64 %SPL, __local_depot8;310; CHECK-NEXT: cvta.local.u64 %SP, %SPL;311; CHECK-NEXT: ld.param.b32 %rd1, [overlapping_chains_param_0];312; CHECK-NEXT: and.b64 %rd2, %rd1, 3;313; CHECK-NEXT: shl.b64 %rd3, %rd2, 2;314; CHECK-NEXT: add.u64 %rd4, %SP, 16;315; CHECK-NEXT: add.s64 %rd5, %rd4, %rd3;316; CHECK-NEXT: st.b32 [%rd5], 10;317; CHECK-NEXT: add.u64 %rd6, %SP, 0;318; CHECK-NEXT: add.s64 %rd7, %rd6, %rd3;319; CHECK-NEXT: ld.b32 %r1, [%SP+28];320; CHECK-NEXT: ld.b32 %r2, [%SP+16];321; CHECK-NEXT: ld.param.b64 %rd8, [overlapping_chains_param_2];322; CHECK-NEXT: st.b32 [%rd7], 10;323; CHECK-NEXT: ld.param.b32 %rd9, [overlapping_chains_param_1];324; CHECK-NEXT: and.b64 %rd10, %rd9, 3;325; CHECK-NEXT: shl.b64 %rd11, %rd10, 2;326; CHECK-NEXT: add.s64 %rd12, %rd6, %rd11;327; CHECK-NEXT: st.b32 [%SP+4], 20;328; CHECK-NEXT: st.b32 [%rd12], 30;329; CHECK-NEXT: ld.param.b64 %rd13, [overlapping_chains_param_3];330; CHECK-NEXT: ld.b32 %r3, [%SP+12];331; CHECK-NEXT: ld.b32 %r4, [%SP+8];332; CHECK-NEXT: ld.b32 %r5, [%SP+4];333; CHECK-NEXT: ld.b32 %r6, [%SP];334; CHECK-NEXT: st.v4.b32 [%rd8], {%r2, 20, 40, %r1};335; CHECK-NEXT: st.v4.b32 [%rd13], {%r6, %r5, %r4, %r3};336; CHECK-NEXT: ret;337 %v0 = insertelement <4 x i32> poison, i32 10, i32 %idx0338 %v1 = insertelement <4 x i32> %v0, i32 20, i32 1339 340 ; Chain 2 starts from v1341 %w0 = insertelement <4 x i32> %v1, i32 30, i32 %idx1342 343 ; Continue chain 1344 %v2 = insertelement <4 x i32> %v1, i32 40, i32 2345 346 store <4 x i32> %v2, ptr %out0347 store <4 x i32> %w0, ptr %out1348 ret void349}350 351; Test with i1 elements (1-bit, non-byte-aligned)352define <8 x i1> @dynamic_i1(i32 %idx) {353; CHECK-LABEL: dynamic_i1(354; CHECK: {355; CHECK-NEXT: .local .align 8 .b8 __local_depot9[8];356; CHECK-NEXT: .reg .b64 %SP;357; CHECK-NEXT: .reg .b64 %SPL;358; CHECK-NEXT: .reg .b32 %r<9>;359; CHECK-NEXT: .reg .b64 %rd<5>;360; CHECK-EMPTY:361; CHECK-NEXT: // %bb.0:362; CHECK-NEXT: mov.b64 %SPL, __local_depot9;363; CHECK-NEXT: cvta.local.u64 %SP, %SPL;364; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i1_param_0];365; CHECK-NEXT: and.b64 %rd2, %rd1, 7;366; CHECK-NEXT: add.u64 %rd3, %SP, 0;367; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;368; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};369; CHECK-NEXT: st.b8 [%rd4], 1;370; CHECK-NEXT: ld.b32 %r3, [%SP];371; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x7773U;372; CHECK-NEXT: ld.b32 %r5, [%SP+4];373; CHECK-NEXT: prmt.b32 %r6, %r5, 0, 0x7771U;374; CHECK-NEXT: prmt.b32 %r7, %r5, 0, 0x7772U;375; CHECK-NEXT: prmt.b32 %r8, %r5, 0, 0x7773U;376; CHECK-NEXT: st.param.b8 [func_retval0+4], %r5;377; CHECK-NEXT: st.param.b8 [func_retval0], %r3;378; CHECK-NEXT: st.param.b8 [func_retval0+7], %r8;379; CHECK-NEXT: st.param.b8 [func_retval0+6], %r7;380; CHECK-NEXT: st.param.b8 [func_retval0+5], %r6;381; CHECK-NEXT: st.param.b8 [func_retval0+3], %r4;382; CHECK-NEXT: st.param.b8 [func_retval0+2], 1;383; CHECK-NEXT: st.param.b8 [func_retval0+1], 0;384; CHECK-NEXT: ret;385 %v0 = insertelement <8 x i1> poison, i1 1, i32 %idx386 %v1 = insertelement <8 x i1> %v0, i1 0, i32 1387 %v2 = insertelement <8 x i1> %v1, i1 1, i32 2388 ret <8 x i1> %v2389}390 391; Test with i2 elements (2-bit, non-byte-aligned)392define <8 x i2> @dynamic_i2(i32 %idx) {393; CHECK-LABEL: dynamic_i2(394; CHECK: {395; CHECK-NEXT: .local .align 8 .b8 __local_depot10[16];396; CHECK-NEXT: .reg .b64 %SP;397; CHECK-NEXT: .reg .b64 %SPL;398; CHECK-NEXT: .reg .b16 %rs<24>;399; CHECK-NEXT: .reg .b32 %r<10>;400; CHECK-NEXT: .reg .b64 %rd<5>;401; CHECK-EMPTY:402; CHECK-NEXT: // %bb.0:403; CHECK-NEXT: mov.b64 %SPL, __local_depot10;404; CHECK-NEXT: cvta.local.u64 %SP, %SPL;405; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i2_param_0];406; CHECK-NEXT: and.b64 %rd2, %rd1, 7;407; CHECK-NEXT: add.u64 %rd3, %SP, 0;408; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;409; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};410; CHECK-NEXT: st.b8 [%rd4], 1;411; CHECK-NEXT: ld.b32 %r3, [%SP+4];412; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;413; CHECK-NEXT: and.b16 %rs2, %rs1, 3;414; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x7771U;415; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;416; CHECK-NEXT: and.b16 %rs4, %rs3, 3;417; CHECK-NEXT: shl.b16 %rs5, %rs4, 2;418; CHECK-NEXT: or.b16 %rs6, %rs2, %rs5;419; CHECK-NEXT: prmt.b32 %r5, %r3, 0, 0x7772U;420; CHECK-NEXT: cvt.u16.u32 %rs7, %r5;421; CHECK-NEXT: and.b16 %rs8, %rs7, 3;422; CHECK-NEXT: shl.b16 %rs9, %rs8, 4;423; CHECK-NEXT: or.b16 %rs10, %rs6, %rs9;424; CHECK-NEXT: prmt.b32 %r6, %r3, 0, 0x7773U;425; CHECK-NEXT: cvt.u16.u32 %rs11, %r6;426; CHECK-NEXT: shl.b16 %rs12, %rs11, 6;427; CHECK-NEXT: or.b16 %rs13, %rs10, %rs12;428; CHECK-NEXT: st.b8 [%SP+8], %rs13;429; CHECK-NEXT: ld.b32 %r7, [%SP];430; CHECK-NEXT: prmt.b32 %r8, %r7, 0, 0x7773U;431; CHECK-NEXT: cvt.u16.u32 %rs14, %r8;432; CHECK-NEXT: shl.b16 %rs15, %rs14, 6;433; CHECK-NEXT: and.b16 %rs16, %rs15, 192;434; CHECK-NEXT: ld.s8 %rs17, [%SP+8];435; CHECK-NEXT: shl.b16 %rs18, %rs17, 8;436; CHECK-NEXT: or.b16 %rs19, %rs16, %rs18;437; CHECK-NEXT: prmt.b32 %r9, %r7, 0, 0x7770U;438; CHECK-NEXT: st.param.b16 [func_retval0], %r9;439; CHECK-NEXT: st.param.b16 [func_retval0+8], %rs17;440; CHECK-NEXT: shr.s16 %rs20, %rs18, 14;441; CHECK-NEXT: st.param.b16 [func_retval0+14], %rs20;442; CHECK-NEXT: shr.s16 %rs21, %rs18, 12;443; CHECK-NEXT: st.param.b16 [func_retval0+12], %rs21;444; CHECK-NEXT: shr.s16 %rs22, %rs18, 10;445; CHECK-NEXT: st.param.b16 [func_retval0+10], %rs22;446; CHECK-NEXT: shr.s16 %rs23, %rs19, 6;447; CHECK-NEXT: st.param.b16 [func_retval0+6], %rs23;448; CHECK-NEXT: st.param.b16 [func_retval0+4], 3;449; CHECK-NEXT: st.param.b16 [func_retval0+2], 2;450; CHECK-NEXT: ret;451 %v0 = insertelement <8 x i2> poison, i2 1, i32 %idx452 %v1 = insertelement <8 x i2> %v0, i2 2, i32 1453 %v2 = insertelement <8 x i2> %v1, i2 3, i32 2454 ret <8 x i2> %v2455}456 457; Test with i3 elements (3-bit, non-byte-aligned)458define <8 x i3> @dynamic_i3(i32 %idx) {459; CHECK-LABEL: dynamic_i3(460; CHECK: {461; CHECK-NEXT: .local .align 8 .b8 __local_depot11[8];462; CHECK-NEXT: .reg .b64 %SP;463; CHECK-NEXT: .reg .b64 %SPL;464; CHECK-NEXT: .reg .b16 %rs<5>;465; CHECK-NEXT: .reg .b32 %r<15>;466; CHECK-NEXT: .reg .b64 %rd<5>;467; CHECK-EMPTY:468; CHECK-NEXT: // %bb.0:469; CHECK-NEXT: mov.b64 %SPL, __local_depot11;470; CHECK-NEXT: cvta.local.u64 %SP, %SPL;471; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i3_param_0];472; CHECK-NEXT: and.b64 %rd2, %rd1, 7;473; CHECK-NEXT: add.u64 %rd3, %SP, 0;474; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;475; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};476; CHECK-NEXT: st.b8 [%rd4], 1;477; CHECK-NEXT: ld.b32 %r3, [%SP];478; CHECK-NEXT: ld.b32 %r4, [%SP+4];479; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;480; CHECK-NEXT: prmt.b32 %r6, %r4, 0, 0x7772U;481; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 0x5410U;482; CHECK-NEXT: st.param.b32 [func_retval0+12], %r7;483; CHECK-NEXT: prmt.b32 %r8, %r4, 0, 0x7771U;484; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7770U;485; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x5410U;486; CHECK-NEXT: st.param.b32 [func_retval0+8], %r10;487; CHECK-NEXT: prmt.b32 %r11, %r3, 0, 0x7773U;488; CHECK-NEXT: cvt.u16.u32 %rs1, %r11;489; CHECK-NEXT: mov.b16 %rs2, 3;490; CHECK-NEXT: mov.b32 %r12, {%rs2, %rs1};491; CHECK-NEXT: st.param.b32 [func_retval0+4], %r12;492; CHECK-NEXT: prmt.b32 %r13, %r3, 0, 0x7770U;493; CHECK-NEXT: cvt.u16.u32 %rs3, %r13;494; CHECK-NEXT: mov.b16 %rs4, 2;495; CHECK-NEXT: mov.b32 %r14, {%rs3, %rs4};496; CHECK-NEXT: st.param.b32 [func_retval0], %r14;497; CHECK-NEXT: ret;498 %v0 = insertelement <8 x i3> poison, i3 1, i32 %idx499 %v1 = insertelement <8 x i3> %v0, i3 2, i32 1500 %v2 = insertelement <8 x i3> %v1, i3 3, i32 2501 ret <8 x i3> %v2502}503 504; Test with i4 elements (4-bit, non-byte-aligned)505define <8 x i4> @dynamic_i4(i32 %idx) {506; CHECK-LABEL: dynamic_i4(507; CHECK: {508; CHECK-NEXT: .local .align 8 .b8 __local_depot12[16];509; CHECK-NEXT: .reg .b64 %SP;510; CHECK-NEXT: .reg .b64 %SPL;511; CHECK-NEXT: .reg .b16 %rs<30>;512; CHECK-NEXT: .reg .b32 %r<22>;513; CHECK-NEXT: .reg .b64 %rd<5>;514; CHECK-EMPTY:515; CHECK-NEXT: // %bb.0:516; CHECK-NEXT: mov.b64 %SPL, __local_depot12;517; CHECK-NEXT: cvta.local.u64 %SP, %SPL;518; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i4_param_0];519; CHECK-NEXT: and.b64 %rd2, %rd1, 7;520; CHECK-NEXT: add.u64 %rd3, %SP, 0;521; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;522; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};523; CHECK-NEXT: st.b8 [%rd4], 1;524; CHECK-NEXT: ld.b32 %r3, [%SP];525; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x7770U;526; CHECK-NEXT: cvt.u16.u32 %rs1, %r4;527; CHECK-NEXT: and.b16 %rs2, %rs1, 15;528; CHECK-NEXT: prmt.b32 %r5, %r3, 0, 0x7771U;529; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;530; CHECK-NEXT: and.b16 %rs4, %rs3, 15;531; CHECK-NEXT: shl.b16 %rs5, %rs4, 4;532; CHECK-NEXT: or.b16 %rs6, %rs2, %rs5;533; CHECK-NEXT: prmt.b32 %r6, %r3, 0, 0x7772U;534; CHECK-NEXT: cvt.u16.u32 %rs7, %r6;535; CHECK-NEXT: and.b16 %rs8, %rs7, 15;536; CHECK-NEXT: shl.b16 %rs9, %rs8, 8;537; CHECK-NEXT: or.b16 %rs10, %rs6, %rs9;538; CHECK-NEXT: prmt.b32 %r7, %r3, 0, 0x7773U;539; CHECK-NEXT: cvt.u16.u32 %rs11, %r7;540; CHECK-NEXT: shl.b16 %rs12, %rs11, 12;541; CHECK-NEXT: or.b16 %rs13, %rs10, %rs12;542; CHECK-NEXT: cvt.u32.u16 %r8, %rs13;543; CHECK-NEXT: ld.b32 %r9, [%SP+4];544; CHECK-NEXT: prmt.b32 %r10, %r9, 0, 0x7770U;545; CHECK-NEXT: cvt.u16.u32 %rs14, %r10;546; CHECK-NEXT: and.b16 %rs15, %rs14, 15;547; CHECK-NEXT: prmt.b32 %r11, %r9, 0, 0x7771U;548; CHECK-NEXT: cvt.u16.u32 %rs16, %r11;549; CHECK-NEXT: and.b16 %rs17, %rs16, 15;550; CHECK-NEXT: shl.b16 %rs18, %rs17, 4;551; CHECK-NEXT: or.b16 %rs19, %rs15, %rs18;552; CHECK-NEXT: prmt.b32 %r12, %r9, 0, 0x7772U;553; CHECK-NEXT: cvt.u16.u32 %rs20, %r12;554; CHECK-NEXT: and.b16 %rs21, %rs20, 15;555; CHECK-NEXT: shl.b16 %rs22, %rs21, 8;556; CHECK-NEXT: or.b16 %rs23, %rs19, %rs22;557; CHECK-NEXT: prmt.b32 %r13, %r9, 0, 0x7773U;558; CHECK-NEXT: cvt.u16.u32 %rs24, %r13;559; CHECK-NEXT: shl.b16 %rs25, %rs24, 12;560; CHECK-NEXT: or.b16 %rs26, %rs23, %rs25;561; CHECK-NEXT: cvt.u32.u16 %r14, %rs26;562; CHECK-NEXT: shl.b32 %r15, %r14, 16;563; CHECK-NEXT: or.b32 %r16, %r8, %r15;564; CHECK-NEXT: mov.b32 %r17, {%rs20, %rs24};565; CHECK-NEXT: st.param.b32 [func_retval0+12], %r17;566; CHECK-NEXT: mov.b32 %r18, {%rs14, %rs16};567; CHECK-NEXT: st.param.b32 [func_retval0+8], %r18;568; CHECK-NEXT: mov.b16 %rs27, 2;569; CHECK-NEXT: mov.b32 %r19, {%rs1, %rs27};570; CHECK-NEXT: st.param.b32 [func_retval0], %r19;571; CHECK-NEXT: shr.u32 %r20, %r16, 12;572; CHECK-NEXT: cvt.u16.u32 %rs28, %r20;573; CHECK-NEXT: mov.b16 %rs29, 3;574; CHECK-NEXT: mov.b32 %r21, {%rs29, %rs28};575; CHECK-NEXT: st.param.b32 [func_retval0+4], %r21;576; CHECK-NEXT: ret;577 %v0 = insertelement <8 x i4> poison, i4 1, i32 %idx578 %v1 = insertelement <8 x i4> %v0, i4 2, i32 1579 %v2 = insertelement <8 x i4> %v1, i4 3, i32 2580 ret <8 x i4> %v2581}582 583; Test with i5 elements (5-bit, non-byte-aligned)584define <8 x i5> @dynamic_i5(i32 %idx) {585; CHECK-LABEL: dynamic_i5(586; CHECK: {587; CHECK-NEXT: .local .align 8 .b8 __local_depot13[8];588; CHECK-NEXT: .reg .b64 %SP;589; CHECK-NEXT: .reg .b64 %SPL;590; CHECK-NEXT: .reg .b16 %rs<5>;591; CHECK-NEXT: .reg .b32 %r<15>;592; CHECK-NEXT: .reg .b64 %rd<5>;593; CHECK-EMPTY:594; CHECK-NEXT: // %bb.0:595; CHECK-NEXT: mov.b64 %SPL, __local_depot13;596; CHECK-NEXT: cvta.local.u64 %SP, %SPL;597; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i5_param_0];598; CHECK-NEXT: and.b64 %rd2, %rd1, 7;599; CHECK-NEXT: add.u64 %rd3, %SP, 0;600; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;601; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};602; CHECK-NEXT: st.b8 [%rd4], 1;603; CHECK-NEXT: ld.b32 %r3, [%SP];604; CHECK-NEXT: ld.b32 %r4, [%SP+4];605; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;606; CHECK-NEXT: prmt.b32 %r6, %r4, 0, 0x7772U;607; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 0x5410U;608; CHECK-NEXT: prmt.b32 %r8, %r4, 0, 0x7771U;609; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7770U;610; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x5410U;611; CHECK-NEXT: st.param.v2.b32 [func_retval0+8], {%r10, %r7};612; CHECK-NEXT: prmt.b32 %r11, %r3, 0, 0x7773U;613; CHECK-NEXT: cvt.u16.u32 %rs1, %r11;614; CHECK-NEXT: mov.b16 %rs2, 3;615; CHECK-NEXT: mov.b32 %r12, {%rs2, %rs1};616; CHECK-NEXT: prmt.b32 %r13, %r3, 0, 0x7770U;617; CHECK-NEXT: cvt.u16.u32 %rs3, %r13;618; CHECK-NEXT: mov.b16 %rs4, 2;619; CHECK-NEXT: mov.b32 %r14, {%rs3, %rs4};620; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r12};621; CHECK-NEXT: ret;622 %v0 = insertelement <8 x i5> poison, i5 1, i32 %idx623 %v1 = insertelement <8 x i5> %v0, i5 2, i32 1624 %v2 = insertelement <8 x i5> %v1, i5 3, i32 2625 ret <8 x i5> %v2626}627 628; Test with i7 elements (7-bit, non-byte-aligned)629define <8 x i7> @dynamic_i7(i32 %idx) {630; CHECK-LABEL: dynamic_i7(631; CHECK: {632; CHECK-NEXT: .local .align 8 .b8 __local_depot14[8];633; CHECK-NEXT: .reg .b64 %SP;634; CHECK-NEXT: .reg .b64 %SPL;635; CHECK-NEXT: .reg .b16 %rs<5>;636; CHECK-NEXT: .reg .b32 %r<15>;637; CHECK-NEXT: .reg .b64 %rd<5>;638; CHECK-EMPTY:639; CHECK-NEXT: // %bb.0:640; CHECK-NEXT: mov.b64 %SPL, __local_depot14;641; CHECK-NEXT: cvta.local.u64 %SP, %SPL;642; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i7_param_0];643; CHECK-NEXT: and.b64 %rd2, %rd1, 7;644; CHECK-NEXT: add.u64 %rd3, %SP, 0;645; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;646; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};647; CHECK-NEXT: st.b8 [%rd4], 1;648; CHECK-NEXT: ld.b32 %r3, [%SP];649; CHECK-NEXT: ld.b32 %r4, [%SP+4];650; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;651; CHECK-NEXT: prmt.b32 %r6, %r4, 0, 0x7772U;652; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 0x5410U;653; CHECK-NEXT: prmt.b32 %r8, %r4, 0, 0x7771U;654; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7770U;655; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x5410U;656; CHECK-NEXT: st.param.v2.b32 [func_retval0+8], {%r10, %r7};657; CHECK-NEXT: prmt.b32 %r11, %r3, 0, 0x7773U;658; CHECK-NEXT: cvt.u16.u32 %rs1, %r11;659; CHECK-NEXT: mov.b16 %rs2, 3;660; CHECK-NEXT: mov.b32 %r12, {%rs2, %rs1};661; CHECK-NEXT: prmt.b32 %r13, %r3, 0, 0x7770U;662; CHECK-NEXT: cvt.u16.u32 %rs3, %r13;663; CHECK-NEXT: mov.b16 %rs4, 2;664; CHECK-NEXT: mov.b32 %r14, {%rs3, %rs4};665; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r12};666; CHECK-NEXT: ret;667 %v0 = insertelement <8 x i7> poison, i7 1, i32 %idx668 %v1 = insertelement <8 x i7> %v0, i7 2, i32 1669 %v2 = insertelement <8 x i7> %v1, i7 3, i32 2670 ret <8 x i7> %v2671}672 673; Test with i6 elements (6-bit, non-byte-aligned)674define <8 x i6> @dynamic_i6(i32 %idx) {675; CHECK-LABEL: dynamic_i6(676; CHECK: {677; CHECK-NEXT: .local .align 8 .b8 __local_depot15[8];678; CHECK-NEXT: .reg .b64 %SP;679; CHECK-NEXT: .reg .b64 %SPL;680; CHECK-NEXT: .reg .b16 %rs<5>;681; CHECK-NEXT: .reg .b32 %r<15>;682; CHECK-NEXT: .reg .b64 %rd<5>;683; CHECK-EMPTY:684; CHECK-NEXT: // %bb.0:685; CHECK-NEXT: mov.b64 %SPL, __local_depot15;686; CHECK-NEXT: cvta.local.u64 %SP, %SPL;687; CHECK-NEXT: ld.param.b32 %rd1, [dynamic_i6_param_0];688; CHECK-NEXT: and.b64 %rd2, %rd1, 7;689; CHECK-NEXT: add.u64 %rd3, %SP, 0;690; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;691; CHECK-NEXT: st.v2.b32 [%SP], {%r1, %r2};692; CHECK-NEXT: st.b8 [%rd4], 1;693; CHECK-NEXT: ld.b32 %r3, [%SP];694; CHECK-NEXT: ld.b32 %r4, [%SP+4];695; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7773U;696; CHECK-NEXT: prmt.b32 %r6, %r4, 0, 0x7772U;697; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 0x5410U;698; CHECK-NEXT: prmt.b32 %r8, %r4, 0, 0x7771U;699; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7770U;700; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x5410U;701; CHECK-NEXT: st.param.v2.b32 [func_retval0+8], {%r10, %r7};702; CHECK-NEXT: prmt.b32 %r11, %r3, 0, 0x7773U;703; CHECK-NEXT: cvt.u16.u32 %rs1, %r11;704; CHECK-NEXT: mov.b16 %rs2, 3;705; CHECK-NEXT: mov.b32 %r12, {%rs2, %rs1};706; CHECK-NEXT: prmt.b32 %r13, %r3, 0, 0x7770U;707; CHECK-NEXT: cvt.u16.u32 %rs3, %r13;708; CHECK-NEXT: mov.b16 %rs4, 2;709; CHECK-NEXT: mov.b32 %r14, {%rs3, %rs4};710; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r12};711; CHECK-NEXT: ret;712 %v0 = insertelement <8 x i6> poison, i6 1, i32 %idx713 %v1 = insertelement <8 x i6> %v0, i6 2, i32 1714 %v2 = insertelement <8 x i6> %v1, i6 3, i32 2715 ret <8 x i6> %v2716}717 718; Test with multiple dynamic insertions on i3 elements719define <4 x i3> @multiple_dynamic_i3(i32 %idx0, i32 %idx1) {720; CHECK-LABEL: multiple_dynamic_i3(721; CHECK: {722; CHECK-NEXT: .reg .b32 %r<9>;723; CHECK-EMPTY:724; CHECK-NEXT: // %bb.0:725; CHECK-NEXT: ld.param.b32 %r1, [multiple_dynamic_i3_param_0];726; CHECK-NEXT: shl.b32 %r2, %r1, 3;727; CHECK-NEXT: bfi.b32 %r3, 1, %r4, %r2, 8;728; CHECK-NEXT: ld.param.b32 %r5, [multiple_dynamic_i3_param_1];729; CHECK-NEXT: shl.b32 %r6, %r5, 3;730; CHECK-NEXT: bfi.b32 %r7, 2, %r3, %r6, 8;731; CHECK-NEXT: st.param.b16 [func_retval0], %r7;732; CHECK-NEXT: shr.u32 %r8, %r7, 16;733; CHECK-NEXT: st.param.b16 [func_retval0+2], %r8;734; CHECK-NEXT: ret;735 %v0 = insertelement <4 x i3> poison, i3 1, i32 %idx0736 %v1 = insertelement <4 x i3> %v0, i3 2, i32 %idx1737 ret <4 x i3> %v1738}739