brintos / llvm-project-archived public Read only

1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -o - -mcpu=sm_100a -march=nvptx64 -mattr=+ptx86 | FileCheck %s3; RUN: llc < %s -o - -mcpu=sm_101a -march=nvptx64 -mattr=+ptx86 | FileCheck %s4; RUN: llc < %s -o - -mcpu=sm_103a -march=nvptx64 -mattr=+ptx88 | FileCheck %s5; RUN: llc < %s -o - -mcpu=sm_100f -march=nvptx64 -mattr=+ptx88 | FileCheck %s6; RUN: llc < %s -o - -mcpu=sm_110f -march=nvptx64 -mattr=+ptx90 | FileCheck %s7; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}8; RUN: %if ptxas-sm_101a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_101a -mattr=+ptx86 | %ptxas-verify -arch=sm_101a %}9; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}10; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}11; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}12 13; CHECK-LABEL: nvvm_tcgen05_st_16x64b14define void @nvvm_tcgen05_st_16x64b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {15; CHECK-LABEL: nvvm_tcgen05_st_16x64b(16; CHECK:       {17; CHECK-NEXT:    .reg .b32 %r<257>;18; CHECK-EMPTY:19; CHECK-NEXT:  // %bb.0:20; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x64b_param_0];21; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_16x64b_param_1];22; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x1.b32 [%r1], {%r2};23; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_16x64b_param_2];24; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x2.b32 [%r1], {%r3, %r4};25; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_16x64b_param_3];26; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x4.b32 [%r1], {%r5, %r6, %r7, %r8};27; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_16x64b_param_4+16];28; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_16x64b_param_4];29; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x8.b32 [%r1], {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};30; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_16x64b_param_5+48];31; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_16x64b_param_5+32];32; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_16x64b_param_5+16];33; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_16x64b_param_5];34; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x16.b32 [%r1], {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};35; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_16x64b_param_6+112];36; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_16x64b_param_6+96];37; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_16x64b_param_6+80];38; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_16x64b_param_6+64];39; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_16x64b_param_6+48];40; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_16x64b_param_6+32];41; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_16x64b_param_6+16];42; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_16x64b_param_6];43; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x32.b32 [%r1], {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};44; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_16x64b_param_7+240];45; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_16x64b_param_7+224];46; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_16x64b_param_7+208];47; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_16x64b_param_7+192];48; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_16x64b_param_7+176];49; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_16x64b_param_7+160];50; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_16x64b_param_7+144];51; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_16x64b_param_7+128];52; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_16x64b_param_7+112];53; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_16x64b_param_7+96];54; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_16x64b_param_7+80];55; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_16x64b_param_7+64];56; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_16x64b_param_7+48];57; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_16x64b_param_7+32];58; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_16x64b_param_7+16];59; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_16x64b_param_7];60; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x64.b32 [%r1], {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};61; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_16x64b_param_8+496];62; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_16x64b_param_8+480];63; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_16x64b_param_8+464];64; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_16x64b_param_8+448];65; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_16x64b_param_8+432];66; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_16x64b_param_8+416];67; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_16x64b_param_8+400];68; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_16x64b_param_8+384];69; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_16x64b_param_8+368];70; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_16x64b_param_8+352];71; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_16x64b_param_8+336];72; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_16x64b_param_8+320];73; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_16x64b_param_8+304];74; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_16x64b_param_8+288];75; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_16x64b_param_8+272];76; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_16x64b_param_8+256];77; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_16x64b_param_8+240];78; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_16x64b_param_8+224];79; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_16x64b_param_8+208];80; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_16x64b_param_8+192];81; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_16x64b_param_8+176];82; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_16x64b_param_8+160];83; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_16x64b_param_8+144];84; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_16x64b_param_8+128];85; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_16x64b_param_8+112];86; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_16x64b_param_8+96];87; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_16x64b_param_8+80];88; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_16x64b_param_8+64];89; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_16x64b_param_8+48];90; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_16x64b_param_8+32];91; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_16x64b_param_8+16];92; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x64b_param_8];93; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x128.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};94; CHECK-NEXT:    ret;95  tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 0)96 97  tail call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 0)98 99  tail call void @llvm.nvvm.tcgen05.st.16x64b.x4(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 0)100 101  tail call void @llvm.nvvm.tcgen05.st.16x64b.x8(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 0)102 103  tail call void @llvm.nvvm.tcgen05.st.16x64b.x16(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 0)104 105  tail call void @llvm.nvvm.tcgen05.st.16x64b.x32(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 0)106 107  tail call void @llvm.nvvm.tcgen05.st.16x64b.x64(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 0)108 109  tail call void @llvm.nvvm.tcgen05.st.16x64b.x128(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 0)110  ret void111}112 113; CHECK-LABEL: nvvm_tcgen05_st_16x64b_unpack114define void @nvvm_tcgen05_st_16x64b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {115; CHECK-LABEL: nvvm_tcgen05_st_16x64b_unpack(116; CHECK:       {117; CHECK-NEXT:    .reg .b32 %r<257>;118; CHECK-EMPTY:119; CHECK-NEXT:  // %bb.0:120; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x64b_unpack_param_0];121; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_16x64b_unpack_param_1];122; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x1.unpack::16b.b32 [%r1], {%r2};123; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_16x64b_unpack_param_2];124; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x2.unpack::16b.b32 [%r1], {%r3, %r4};125; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_16x64b_unpack_param_3];126; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x4.unpack::16b.b32 [%r1], {%r5, %r6, %r7, %r8};127; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_16x64b_unpack_param_4+16];128; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_16x64b_unpack_param_4];129; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x8.unpack::16b.b32 [%r1], {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};130; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_16x64b_unpack_param_5+48];131; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_16x64b_unpack_param_5+32];132; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_16x64b_unpack_param_5+16];133; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_16x64b_unpack_param_5];134; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x16.unpack::16b.b32 [%r1], {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};135; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_16x64b_unpack_param_6+112];136; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_16x64b_unpack_param_6+96];137; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_16x64b_unpack_param_6+80];138; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_16x64b_unpack_param_6+64];139; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_16x64b_unpack_param_6+48];140; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_16x64b_unpack_param_6+32];141; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_16x64b_unpack_param_6+16];142; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_16x64b_unpack_param_6];143; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x32.unpack::16b.b32 [%r1], {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};144; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_16x64b_unpack_param_7+240];145; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_16x64b_unpack_param_7+224];146; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_16x64b_unpack_param_7+208];147; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_16x64b_unpack_param_7+192];148; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_16x64b_unpack_param_7+176];149; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_16x64b_unpack_param_7+160];150; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_16x64b_unpack_param_7+144];151; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_16x64b_unpack_param_7+128];152; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_16x64b_unpack_param_7+112];153; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_16x64b_unpack_param_7+96];154; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_16x64b_unpack_param_7+80];155; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_16x64b_unpack_param_7+64];156; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_16x64b_unpack_param_7+48];157; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_16x64b_unpack_param_7+32];158; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_16x64b_unpack_param_7+16];159; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_16x64b_unpack_param_7];160; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x64.unpack::16b.b32 [%r1], {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};161; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_16x64b_unpack_param_8+496];162; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_16x64b_unpack_param_8+480];163; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_16x64b_unpack_param_8+464];164; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_16x64b_unpack_param_8+448];165; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_16x64b_unpack_param_8+432];166; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_16x64b_unpack_param_8+416];167; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_16x64b_unpack_param_8+400];168; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_16x64b_unpack_param_8+384];169; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_16x64b_unpack_param_8+368];170; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_16x64b_unpack_param_8+352];171; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_16x64b_unpack_param_8+336];172; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_16x64b_unpack_param_8+320];173; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_16x64b_unpack_param_8+304];174; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_16x64b_unpack_param_8+288];175; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_16x64b_unpack_param_8+272];176; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_16x64b_unpack_param_8+256];177; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_16x64b_unpack_param_8+240];178; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_16x64b_unpack_param_8+224];179; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_16x64b_unpack_param_8+208];180; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_16x64b_unpack_param_8+192];181; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_16x64b_unpack_param_8+176];182; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_16x64b_unpack_param_8+160];183; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_16x64b_unpack_param_8+144];184; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_16x64b_unpack_param_8+128];185; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_16x64b_unpack_param_8+112];186; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_16x64b_unpack_param_8+96];187; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_16x64b_unpack_param_8+80];188; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_16x64b_unpack_param_8+64];189; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_16x64b_unpack_param_8+48];190; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_16x64b_unpack_param_8+32];191; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_16x64b_unpack_param_8+16];192; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x64b_unpack_param_8];193; CHECK-NEXT:    tcgen05.st.sync.aligned.16x64b.x128.unpack::16b.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};194; CHECK-NEXT:    ret;195  tail call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 1)196 197  tail call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 1)198 199  tail call void @llvm.nvvm.tcgen05.st.16x64b.x4(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 1)200 201  tail call void @llvm.nvvm.tcgen05.st.16x64b.x8(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 1)202 203  tail call void @llvm.nvvm.tcgen05.st.16x64b.x16(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 1)204 205  tail call void @llvm.nvvm.tcgen05.st.16x64b.x32(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 1)206 207  tail call void @llvm.nvvm.tcgen05.st.16x64b.x64(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 1)208 209  tail call void @llvm.nvvm.tcgen05.st.16x64b.x128(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 1)210  ret void211}212 213; CHECK-LABEL: nvvm_tcgen05_st_16x128b214define void @nvvm_tcgen05_st_16x128b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {215; CHECK-LABEL: nvvm_tcgen05_st_16x128b(216; CHECK:       {217; CHECK-NEXT:    .reg .b32 %r<256>;218; CHECK-EMPTY:219; CHECK-NEXT:  // %bb.0:220; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x128b_param_0];221; CHECK-NEXT:    ld.param.v2.b32 {%r2, %r3}, [nvvm_tcgen05_st_16x128b_param_2];222; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x1.b32 [%r1], {%r2, %r3};223; CHECK-NEXT:    ld.param.v4.b32 {%r4, %r5, %r6, %r7}, [nvvm_tcgen05_st_16x128b_param_3];224; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x2.b32 [%r1], {%r4, %r5, %r6, %r7};225; CHECK-NEXT:    ld.param.v4.b32 {%r8, %r9, %r10, %r11}, [nvvm_tcgen05_st_16x128b_param_4+16];226; CHECK-NEXT:    ld.param.v4.b32 {%r12, %r13, %r14, %r15}, [nvvm_tcgen05_st_16x128b_param_4];227; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x4.b32 [%r1], {%r12, %r13, %r14, %r15, %r8, %r9, %r10, %r11};228; CHECK-NEXT:    ld.param.v4.b32 {%r16, %r17, %r18, %r19}, [nvvm_tcgen05_st_16x128b_param_5+48];229; CHECK-NEXT:    ld.param.v4.b32 {%r20, %r21, %r22, %r23}, [nvvm_tcgen05_st_16x128b_param_5+32];230; CHECK-NEXT:    ld.param.v4.b32 {%r24, %r25, %r26, %r27}, [nvvm_tcgen05_st_16x128b_param_5+16];231; CHECK-NEXT:    ld.param.v4.b32 {%r28, %r29, %r30, %r31}, [nvvm_tcgen05_st_16x128b_param_5];232; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x8.b32 [%r1], {%r28, %r29, %r30, %r31, %r24, %r25, %r26, %r27, %r20, %r21, %r22, %r23, %r16, %r17, %r18, %r19};233; CHECK-NEXT:    ld.param.v4.b32 {%r32, %r33, %r34, %r35}, [nvvm_tcgen05_st_16x128b_param_6+112];234; CHECK-NEXT:    ld.param.v4.b32 {%r36, %r37, %r38, %r39}, [nvvm_tcgen05_st_16x128b_param_6+96];235; CHECK-NEXT:    ld.param.v4.b32 {%r40, %r41, %r42, %r43}, [nvvm_tcgen05_st_16x128b_param_6+80];236; CHECK-NEXT:    ld.param.v4.b32 {%r44, %r45, %r46, %r47}, [nvvm_tcgen05_st_16x128b_param_6+64];237; CHECK-NEXT:    ld.param.v4.b32 {%r48, %r49, %r50, %r51}, [nvvm_tcgen05_st_16x128b_param_6+48];238; CHECK-NEXT:    ld.param.v4.b32 {%r52, %r53, %r54, %r55}, [nvvm_tcgen05_st_16x128b_param_6+32];239; CHECK-NEXT:    ld.param.v4.b32 {%r56, %r57, %r58, %r59}, [nvvm_tcgen05_st_16x128b_param_6+16];240; CHECK-NEXT:    ld.param.v4.b32 {%r60, %r61, %r62, %r63}, [nvvm_tcgen05_st_16x128b_param_6];241; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x16.b32 [%r1], {%r60, %r61, %r62, %r63, %r56, %r57, %r58, %r59, %r52, %r53, %r54, %r55, %r48, %r49, %r50, %r51, %r44, %r45, %r46, %r47, %r40, %r41, %r42, %r43, %r36, %r37, %r38, %r39, %r32, %r33, %r34, %r35};242; CHECK-NEXT:    ld.param.v4.b32 {%r64, %r65, %r66, %r67}, [nvvm_tcgen05_st_16x128b_param_7+240];243; CHECK-NEXT:    ld.param.v4.b32 {%r68, %r69, %r70, %r71}, [nvvm_tcgen05_st_16x128b_param_7+224];244; CHECK-NEXT:    ld.param.v4.b32 {%r72, %r73, %r74, %r75}, [nvvm_tcgen05_st_16x128b_param_7+208];245; CHECK-NEXT:    ld.param.v4.b32 {%r76, %r77, %r78, %r79}, [nvvm_tcgen05_st_16x128b_param_7+192];246; CHECK-NEXT:    ld.param.v4.b32 {%r80, %r81, %r82, %r83}, [nvvm_tcgen05_st_16x128b_param_7+176];247; CHECK-NEXT:    ld.param.v4.b32 {%r84, %r85, %r86, %r87}, [nvvm_tcgen05_st_16x128b_param_7+160];248; CHECK-NEXT:    ld.param.v4.b32 {%r88, %r89, %r90, %r91}, [nvvm_tcgen05_st_16x128b_param_7+144];249; CHECK-NEXT:    ld.param.v4.b32 {%r92, %r93, %r94, %r95}, [nvvm_tcgen05_st_16x128b_param_7+128];250; CHECK-NEXT:    ld.param.v4.b32 {%r96, %r97, %r98, %r99}, [nvvm_tcgen05_st_16x128b_param_7+112];251; CHECK-NEXT:    ld.param.v4.b32 {%r100, %r101, %r102, %r103}, [nvvm_tcgen05_st_16x128b_param_7+96];252; CHECK-NEXT:    ld.param.v4.b32 {%r104, %r105, %r106, %r107}, [nvvm_tcgen05_st_16x128b_param_7+80];253; CHECK-NEXT:    ld.param.v4.b32 {%r108, %r109, %r110, %r111}, [nvvm_tcgen05_st_16x128b_param_7+64];254; CHECK-NEXT:    ld.param.v4.b32 {%r112, %r113, %r114, %r115}, [nvvm_tcgen05_st_16x128b_param_7+48];255; CHECK-NEXT:    ld.param.v4.b32 {%r116, %r117, %r118, %r119}, [nvvm_tcgen05_st_16x128b_param_7+32];256; CHECK-NEXT:    ld.param.v4.b32 {%r120, %r121, %r122, %r123}, [nvvm_tcgen05_st_16x128b_param_7+16];257; CHECK-NEXT:    ld.param.v4.b32 {%r124, %r125, %r126, %r127}, [nvvm_tcgen05_st_16x128b_param_7];258; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x32.b32 [%r1], {%r124, %r125, %r126, %r127, %r120, %r121, %r122, %r123, %r116, %r117, %r118, %r119, %r112, %r113, %r114, %r115, %r108, %r109, %r110, %r111, %r104, %r105, %r106, %r107, %r100, %r101, %r102, %r103, %r96, %r97, %r98, %r99, %r92, %r93, %r94, %r95, %r88, %r89, %r90, %r91, %r84, %r85, %r86, %r87, %r80, %r81, %r82, %r83, %r76, %r77, %r78, %r79, %r72, %r73, %r74, %r75, %r68, %r69, %r70, %r71, %r64, %r65, %r66, %r67};259; CHECK-NEXT:    ld.param.v4.b32 {%r128, %r129, %r130, %r131}, [nvvm_tcgen05_st_16x128b_param_8+496];260; CHECK-NEXT:    ld.param.v4.b32 {%r132, %r133, %r134, %r135}, [nvvm_tcgen05_st_16x128b_param_8+480];261; CHECK-NEXT:    ld.param.v4.b32 {%r136, %r137, %r138, %r139}, [nvvm_tcgen05_st_16x128b_param_8+464];262; CHECK-NEXT:    ld.param.v4.b32 {%r140, %r141, %r142, %r143}, [nvvm_tcgen05_st_16x128b_param_8+448];263; CHECK-NEXT:    ld.param.v4.b32 {%r144, %r145, %r146, %r147}, [nvvm_tcgen05_st_16x128b_param_8+432];264; CHECK-NEXT:    ld.param.v4.b32 {%r148, %r149, %r150, %r151}, [nvvm_tcgen05_st_16x128b_param_8+416];265; CHECK-NEXT:    ld.param.v4.b32 {%r152, %r153, %r154, %r155}, [nvvm_tcgen05_st_16x128b_param_8+400];266; CHECK-NEXT:    ld.param.v4.b32 {%r156, %r157, %r158, %r159}, [nvvm_tcgen05_st_16x128b_param_8+384];267; CHECK-NEXT:    ld.param.v4.b32 {%r160, %r161, %r162, %r163}, [nvvm_tcgen05_st_16x128b_param_8+368];268; CHECK-NEXT:    ld.param.v4.b32 {%r164, %r165, %r166, %r167}, [nvvm_tcgen05_st_16x128b_param_8+352];269; CHECK-NEXT:    ld.param.v4.b32 {%r168, %r169, %r170, %r171}, [nvvm_tcgen05_st_16x128b_param_8+336];270; CHECK-NEXT:    ld.param.v4.b32 {%r172, %r173, %r174, %r175}, [nvvm_tcgen05_st_16x128b_param_8+320];271; CHECK-NEXT:    ld.param.v4.b32 {%r176, %r177, %r178, %r179}, [nvvm_tcgen05_st_16x128b_param_8+304];272; CHECK-NEXT:    ld.param.v4.b32 {%r180, %r181, %r182, %r183}, [nvvm_tcgen05_st_16x128b_param_8+288];273; CHECK-NEXT:    ld.param.v4.b32 {%r184, %r185, %r186, %r187}, [nvvm_tcgen05_st_16x128b_param_8+272];274; CHECK-NEXT:    ld.param.v4.b32 {%r188, %r189, %r190, %r191}, [nvvm_tcgen05_st_16x128b_param_8+256];275; CHECK-NEXT:    ld.param.v4.b32 {%r192, %r193, %r194, %r195}, [nvvm_tcgen05_st_16x128b_param_8+240];276; CHECK-NEXT:    ld.param.v4.b32 {%r196, %r197, %r198, %r199}, [nvvm_tcgen05_st_16x128b_param_8+224];277; CHECK-NEXT:    ld.param.v4.b32 {%r200, %r201, %r202, %r203}, [nvvm_tcgen05_st_16x128b_param_8+208];278; CHECK-NEXT:    ld.param.v4.b32 {%r204, %r205, %r206, %r207}, [nvvm_tcgen05_st_16x128b_param_8+192];279; CHECK-NEXT:    ld.param.v4.b32 {%r208, %r209, %r210, %r211}, [nvvm_tcgen05_st_16x128b_param_8+176];280; CHECK-NEXT:    ld.param.v4.b32 {%r212, %r213, %r214, %r215}, [nvvm_tcgen05_st_16x128b_param_8+160];281; CHECK-NEXT:    ld.param.v4.b32 {%r216, %r217, %r218, %r219}, [nvvm_tcgen05_st_16x128b_param_8+144];282; CHECK-NEXT:    ld.param.v4.b32 {%r220, %r221, %r222, %r223}, [nvvm_tcgen05_st_16x128b_param_8+128];283; CHECK-NEXT:    ld.param.v4.b32 {%r224, %r225, %r226, %r227}, [nvvm_tcgen05_st_16x128b_param_8+112];284; CHECK-NEXT:    ld.param.v4.b32 {%r228, %r229, %r230, %r231}, [nvvm_tcgen05_st_16x128b_param_8+96];285; CHECK-NEXT:    ld.param.v4.b32 {%r232, %r233, %r234, %r235}, [nvvm_tcgen05_st_16x128b_param_8+80];286; CHECK-NEXT:    ld.param.v4.b32 {%r236, %r237, %r238, %r239}, [nvvm_tcgen05_st_16x128b_param_8+64];287; CHECK-NEXT:    ld.param.v4.b32 {%r240, %r241, %r242, %r243}, [nvvm_tcgen05_st_16x128b_param_8+48];288; CHECK-NEXT:    ld.param.v4.b32 {%r244, %r245, %r246, %r247}, [nvvm_tcgen05_st_16x128b_param_8+32];289; CHECK-NEXT:    ld.param.v4.b32 {%r248, %r249, %r250, %r251}, [nvvm_tcgen05_st_16x128b_param_8+16];290; CHECK-NEXT:    ld.param.v4.b32 {%r252, %r253, %r254, %r255}, [nvvm_tcgen05_st_16x128b_param_8];291; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x64.b32 [%r1], {%r252, %r253, %r254, %r255, %r248, %r249, %r250, %r251, %r244, %r245, %r246, %r247, %r240, %r241, %r242, %r243, %r236, %r237, %r238, %r239, %r232, %r233, %r234, %r235, %r228, %r229, %r230, %r231, %r224, %r225, %r226, %r227, %r220, %r221, %r222, %r223, %r216, %r217, %r218, %r219, %r212, %r213, %r214, %r215, %r208, %r209, %r210, %r211, %r204, %r205, %r206, %r207, %r200, %r201, %r202, %r203, %r196, %r197, %r198, %r199, %r192, %r193, %r194, %r195, %r188, %r189, %r190, %r191, %r184, %r185, %r186, %r187, %r180, %r181, %r182, %r183, %r176, %r177, %r178, %r179, %r172, %r173, %r174, %r175, %r168, %r169, %r170, %r171, %r164, %r165, %r166, %r167, %r160, %r161, %r162, %r163, %r156, %r157, %r158, %r159, %r152, %r153, %r154, %r155, %r148, %r149, %r150, %r151, %r144, %r145, %r146, %r147, %r140, %r141, %r142, %r143, %r136, %r137, %r138, %r139, %r132, %r133, %r134, %r135, %r128, %r129, %r130, %r131};292; CHECK-NEXT:    ret;293  tail call void @llvm.nvvm.tcgen05.st.16x128b.x1(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 0)294 295  tail call void @llvm.nvvm.tcgen05.st.16x128b.x2(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 0)296 297  tail call void @llvm.nvvm.tcgen05.st.16x128b.x4(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 0)298 299  tail call void @llvm.nvvm.tcgen05.st.16x128b.x8(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 0)300 301  tail call void @llvm.nvvm.tcgen05.st.16x128b.x16(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 0)302 303  tail call void @llvm.nvvm.tcgen05.st.16x128b.x32(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 0)304 305  tail call void @llvm.nvvm.tcgen05.st.16x128b.x64(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 0)306  ret void307}308 309; CHECK-LABEL: nvvm_tcgen05_st_16x128b_unpack310define void @nvvm_tcgen05_st_16x128b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {311; CHECK-LABEL: nvvm_tcgen05_st_16x128b_unpack(312; CHECK:       {313; CHECK-NEXT:    .reg .b32 %r<256>;314; CHECK-EMPTY:315; CHECK-NEXT:  // %bb.0:316; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x128b_unpack_param_0];317; CHECK-NEXT:    ld.param.v2.b32 {%r2, %r3}, [nvvm_tcgen05_st_16x128b_unpack_param_2];318; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x1.unpack::16b.b32 [%r1], {%r2, %r3};319; CHECK-NEXT:    ld.param.v4.b32 {%r4, %r5, %r6, %r7}, [nvvm_tcgen05_st_16x128b_unpack_param_3];320; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x2.unpack::16b.b32 [%r1], {%r4, %r5, %r6, %r7};321; CHECK-NEXT:    ld.param.v4.b32 {%r8, %r9, %r10, %r11}, [nvvm_tcgen05_st_16x128b_unpack_param_4+16];322; CHECK-NEXT:    ld.param.v4.b32 {%r12, %r13, %r14, %r15}, [nvvm_tcgen05_st_16x128b_unpack_param_4];323; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x4.unpack::16b.b32 [%r1], {%r12, %r13, %r14, %r15, %r8, %r9, %r10, %r11};324; CHECK-NEXT:    ld.param.v4.b32 {%r16, %r17, %r18, %r19}, [nvvm_tcgen05_st_16x128b_unpack_param_5+48];325; CHECK-NEXT:    ld.param.v4.b32 {%r20, %r21, %r22, %r23}, [nvvm_tcgen05_st_16x128b_unpack_param_5+32];326; CHECK-NEXT:    ld.param.v4.b32 {%r24, %r25, %r26, %r27}, [nvvm_tcgen05_st_16x128b_unpack_param_5+16];327; CHECK-NEXT:    ld.param.v4.b32 {%r28, %r29, %r30, %r31}, [nvvm_tcgen05_st_16x128b_unpack_param_5];328; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x8.unpack::16b.b32 [%r1], {%r28, %r29, %r30, %r31, %r24, %r25, %r26, %r27, %r20, %r21, %r22, %r23, %r16, %r17, %r18, %r19};329; CHECK-NEXT:    ld.param.v4.b32 {%r32, %r33, %r34, %r35}, [nvvm_tcgen05_st_16x128b_unpack_param_6+112];330; CHECK-NEXT:    ld.param.v4.b32 {%r36, %r37, %r38, %r39}, [nvvm_tcgen05_st_16x128b_unpack_param_6+96];331; CHECK-NEXT:    ld.param.v4.b32 {%r40, %r41, %r42, %r43}, [nvvm_tcgen05_st_16x128b_unpack_param_6+80];332; CHECK-NEXT:    ld.param.v4.b32 {%r44, %r45, %r46, %r47}, [nvvm_tcgen05_st_16x128b_unpack_param_6+64];333; CHECK-NEXT:    ld.param.v4.b32 {%r48, %r49, %r50, %r51}, [nvvm_tcgen05_st_16x128b_unpack_param_6+48];334; CHECK-NEXT:    ld.param.v4.b32 {%r52, %r53, %r54, %r55}, [nvvm_tcgen05_st_16x128b_unpack_param_6+32];335; CHECK-NEXT:    ld.param.v4.b32 {%r56, %r57, %r58, %r59}, [nvvm_tcgen05_st_16x128b_unpack_param_6+16];336; CHECK-NEXT:    ld.param.v4.b32 {%r60, %r61, %r62, %r63}, [nvvm_tcgen05_st_16x128b_unpack_param_6];337; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x16.unpack::16b.b32 [%r1], {%r60, %r61, %r62, %r63, %r56, %r57, %r58, %r59, %r52, %r53, %r54, %r55, %r48, %r49, %r50, %r51, %r44, %r45, %r46, %r47, %r40, %r41, %r42, %r43, %r36, %r37, %r38, %r39, %r32, %r33, %r34, %r35};338; CHECK-NEXT:    ld.param.v4.b32 {%r64, %r65, %r66, %r67}, [nvvm_tcgen05_st_16x128b_unpack_param_7+240];339; CHECK-NEXT:    ld.param.v4.b32 {%r68, %r69, %r70, %r71}, [nvvm_tcgen05_st_16x128b_unpack_param_7+224];340; CHECK-NEXT:    ld.param.v4.b32 {%r72, %r73, %r74, %r75}, [nvvm_tcgen05_st_16x128b_unpack_param_7+208];341; CHECK-NEXT:    ld.param.v4.b32 {%r76, %r77, %r78, %r79}, [nvvm_tcgen05_st_16x128b_unpack_param_7+192];342; CHECK-NEXT:    ld.param.v4.b32 {%r80, %r81, %r82, %r83}, [nvvm_tcgen05_st_16x128b_unpack_param_7+176];343; CHECK-NEXT:    ld.param.v4.b32 {%r84, %r85, %r86, %r87}, [nvvm_tcgen05_st_16x128b_unpack_param_7+160];344; CHECK-NEXT:    ld.param.v4.b32 {%r88, %r89, %r90, %r91}, [nvvm_tcgen05_st_16x128b_unpack_param_7+144];345; CHECK-NEXT:    ld.param.v4.b32 {%r92, %r93, %r94, %r95}, [nvvm_tcgen05_st_16x128b_unpack_param_7+128];346; CHECK-NEXT:    ld.param.v4.b32 {%r96, %r97, %r98, %r99}, [nvvm_tcgen05_st_16x128b_unpack_param_7+112];347; CHECK-NEXT:    ld.param.v4.b32 {%r100, %r101, %r102, %r103}, [nvvm_tcgen05_st_16x128b_unpack_param_7+96];348; CHECK-NEXT:    ld.param.v4.b32 {%r104, %r105, %r106, %r107}, [nvvm_tcgen05_st_16x128b_unpack_param_7+80];349; CHECK-NEXT:    ld.param.v4.b32 {%r108, %r109, %r110, %r111}, [nvvm_tcgen05_st_16x128b_unpack_param_7+64];350; CHECK-NEXT:    ld.param.v4.b32 {%r112, %r113, %r114, %r115}, [nvvm_tcgen05_st_16x128b_unpack_param_7+48];351; CHECK-NEXT:    ld.param.v4.b32 {%r116, %r117, %r118, %r119}, [nvvm_tcgen05_st_16x128b_unpack_param_7+32];352; CHECK-NEXT:    ld.param.v4.b32 {%r120, %r121, %r122, %r123}, [nvvm_tcgen05_st_16x128b_unpack_param_7+16];353; CHECK-NEXT:    ld.param.v4.b32 {%r124, %r125, %r126, %r127}, [nvvm_tcgen05_st_16x128b_unpack_param_7];354; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x32.unpack::16b.b32 [%r1], {%r124, %r125, %r126, %r127, %r120, %r121, %r122, %r123, %r116, %r117, %r118, %r119, %r112, %r113, %r114, %r115, %r108, %r109, %r110, %r111, %r104, %r105, %r106, %r107, %r100, %r101, %r102, %r103, %r96, %r97, %r98, %r99, %r92, %r93, %r94, %r95, %r88, %r89, %r90, %r91, %r84, %r85, %r86, %r87, %r80, %r81, %r82, %r83, %r76, %r77, %r78, %r79, %r72, %r73, %r74, %r75, %r68, %r69, %r70, %r71, %r64, %r65, %r66, %r67};355; CHECK-NEXT:    ld.param.v4.b32 {%r128, %r129, %r130, %r131}, [nvvm_tcgen05_st_16x128b_unpack_param_8+496];356; CHECK-NEXT:    ld.param.v4.b32 {%r132, %r133, %r134, %r135}, [nvvm_tcgen05_st_16x128b_unpack_param_8+480];357; CHECK-NEXT:    ld.param.v4.b32 {%r136, %r137, %r138, %r139}, [nvvm_tcgen05_st_16x128b_unpack_param_8+464];358; CHECK-NEXT:    ld.param.v4.b32 {%r140, %r141, %r142, %r143}, [nvvm_tcgen05_st_16x128b_unpack_param_8+448];359; CHECK-NEXT:    ld.param.v4.b32 {%r144, %r145, %r146, %r147}, [nvvm_tcgen05_st_16x128b_unpack_param_8+432];360; CHECK-NEXT:    ld.param.v4.b32 {%r148, %r149, %r150, %r151}, [nvvm_tcgen05_st_16x128b_unpack_param_8+416];361; CHECK-NEXT:    ld.param.v4.b32 {%r152, %r153, %r154, %r155}, [nvvm_tcgen05_st_16x128b_unpack_param_8+400];362; CHECK-NEXT:    ld.param.v4.b32 {%r156, %r157, %r158, %r159}, [nvvm_tcgen05_st_16x128b_unpack_param_8+384];363; CHECK-NEXT:    ld.param.v4.b32 {%r160, %r161, %r162, %r163}, [nvvm_tcgen05_st_16x128b_unpack_param_8+368];364; CHECK-NEXT:    ld.param.v4.b32 {%r164, %r165, %r166, %r167}, [nvvm_tcgen05_st_16x128b_unpack_param_8+352];365; CHECK-NEXT:    ld.param.v4.b32 {%r168, %r169, %r170, %r171}, [nvvm_tcgen05_st_16x128b_unpack_param_8+336];366; CHECK-NEXT:    ld.param.v4.b32 {%r172, %r173, %r174, %r175}, [nvvm_tcgen05_st_16x128b_unpack_param_8+320];367; CHECK-NEXT:    ld.param.v4.b32 {%r176, %r177, %r178, %r179}, [nvvm_tcgen05_st_16x128b_unpack_param_8+304];368; CHECK-NEXT:    ld.param.v4.b32 {%r180, %r181, %r182, %r183}, [nvvm_tcgen05_st_16x128b_unpack_param_8+288];369; CHECK-NEXT:    ld.param.v4.b32 {%r184, %r185, %r186, %r187}, [nvvm_tcgen05_st_16x128b_unpack_param_8+272];370; CHECK-NEXT:    ld.param.v4.b32 {%r188, %r189, %r190, %r191}, [nvvm_tcgen05_st_16x128b_unpack_param_8+256];371; CHECK-NEXT:    ld.param.v4.b32 {%r192, %r193, %r194, %r195}, [nvvm_tcgen05_st_16x128b_unpack_param_8+240];372; CHECK-NEXT:    ld.param.v4.b32 {%r196, %r197, %r198, %r199}, [nvvm_tcgen05_st_16x128b_unpack_param_8+224];373; CHECK-NEXT:    ld.param.v4.b32 {%r200, %r201, %r202, %r203}, [nvvm_tcgen05_st_16x128b_unpack_param_8+208];374; CHECK-NEXT:    ld.param.v4.b32 {%r204, %r205, %r206, %r207}, [nvvm_tcgen05_st_16x128b_unpack_param_8+192];375; CHECK-NEXT:    ld.param.v4.b32 {%r208, %r209, %r210, %r211}, [nvvm_tcgen05_st_16x128b_unpack_param_8+176];376; CHECK-NEXT:    ld.param.v4.b32 {%r212, %r213, %r214, %r215}, [nvvm_tcgen05_st_16x128b_unpack_param_8+160];377; CHECK-NEXT:    ld.param.v4.b32 {%r216, %r217, %r218, %r219}, [nvvm_tcgen05_st_16x128b_unpack_param_8+144];378; CHECK-NEXT:    ld.param.v4.b32 {%r220, %r221, %r222, %r223}, [nvvm_tcgen05_st_16x128b_unpack_param_8+128];379; CHECK-NEXT:    ld.param.v4.b32 {%r224, %r225, %r226, %r227}, [nvvm_tcgen05_st_16x128b_unpack_param_8+112];380; CHECK-NEXT:    ld.param.v4.b32 {%r228, %r229, %r230, %r231}, [nvvm_tcgen05_st_16x128b_unpack_param_8+96];381; CHECK-NEXT:    ld.param.v4.b32 {%r232, %r233, %r234, %r235}, [nvvm_tcgen05_st_16x128b_unpack_param_8+80];382; CHECK-NEXT:    ld.param.v4.b32 {%r236, %r237, %r238, %r239}, [nvvm_tcgen05_st_16x128b_unpack_param_8+64];383; CHECK-NEXT:    ld.param.v4.b32 {%r240, %r241, %r242, %r243}, [nvvm_tcgen05_st_16x128b_unpack_param_8+48];384; CHECK-NEXT:    ld.param.v4.b32 {%r244, %r245, %r246, %r247}, [nvvm_tcgen05_st_16x128b_unpack_param_8+32];385; CHECK-NEXT:    ld.param.v4.b32 {%r248, %r249, %r250, %r251}, [nvvm_tcgen05_st_16x128b_unpack_param_8+16];386; CHECK-NEXT:    ld.param.v4.b32 {%r252, %r253, %r254, %r255}, [nvvm_tcgen05_st_16x128b_unpack_param_8];387; CHECK-NEXT:    tcgen05.st.sync.aligned.16x128b.x64.unpack::16b.b32 [%r1], {%r252, %r253, %r254, %r255, %r248, %r249, %r250, %r251, %r244, %r245, %r246, %r247, %r240, %r241, %r242, %r243, %r236, %r237, %r238, %r239, %r232, %r233, %r234, %r235, %r228, %r229, %r230, %r231, %r224, %r225, %r226, %r227, %r220, %r221, %r222, %r223, %r216, %r217, %r218, %r219, %r212, %r213, %r214, %r215, %r208, %r209, %r210, %r211, %r204, %r205, %r206, %r207, %r200, %r201, %r202, %r203, %r196, %r197, %r198, %r199, %r192, %r193, %r194, %r195, %r188, %r189, %r190, %r191, %r184, %r185, %r186, %r187, %r180, %r181, %r182, %r183, %r176, %r177, %r178, %r179, %r172, %r173, %r174, %r175, %r168, %r169, %r170, %r171, %r164, %r165, %r166, %r167, %r160, %r161, %r162, %r163, %r156, %r157, %r158, %r159, %r152, %r153, %r154, %r155, %r148, %r149, %r150, %r151, %r144, %r145, %r146, %r147, %r140, %r141, %r142, %r143, %r136, %r137, %r138, %r139, %r132, %r133, %r134, %r135, %r128, %r129, %r130, %r131};388; CHECK-NEXT:    ret;389  tail call void @llvm.nvvm.tcgen05.st.16x128b.x1(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 1)390 391  tail call void @llvm.nvvm.tcgen05.st.16x128b.x2(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 1)392 393  tail call void @llvm.nvvm.tcgen05.st.16x128b.x4(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 1)394 395  tail call void @llvm.nvvm.tcgen05.st.16x128b.x8(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 1)396 397  tail call void @llvm.nvvm.tcgen05.st.16x128b.x16(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 1)398 399  tail call void @llvm.nvvm.tcgen05.st.16x128b.x32(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 1)400 401  tail call void @llvm.nvvm.tcgen05.st.16x128b.x64(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 1)402  ret void403}404 405; CHECK-LABEL: nvvm_tcgen05_st_16x256b406define void @nvvm_tcgen05_st_16x256b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {407; CHECK-LABEL: nvvm_tcgen05_st_16x256b(408; CHECK:       {409; CHECK-NEXT:    .reg .b32 %r<254>;410; CHECK-EMPTY:411; CHECK-NEXT:  // %bb.0:412; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x256b_param_0];413; CHECK-NEXT:    ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [nvvm_tcgen05_st_16x256b_param_3];414; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x1.b32 [%r1], {%r2, %r3, %r4, %r5};415; CHECK-NEXT:    ld.param.v4.b32 {%r6, %r7, %r8, %r9}, [nvvm_tcgen05_st_16x256b_param_4+16];416; CHECK-NEXT:    ld.param.v4.b32 {%r10, %r11, %r12, %r13}, [nvvm_tcgen05_st_16x256b_param_4];417; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x2.b32 [%r1], {%r10, %r11, %r12, %r13, %r6, %r7, %r8, %r9};418; CHECK-NEXT:    ld.param.v4.b32 {%r14, %r15, %r16, %r17}, [nvvm_tcgen05_st_16x256b_param_5+48];419; CHECK-NEXT:    ld.param.v4.b32 {%r18, %r19, %r20, %r21}, [nvvm_tcgen05_st_16x256b_param_5+32];420; CHECK-NEXT:    ld.param.v4.b32 {%r22, %r23, %r24, %r25}, [nvvm_tcgen05_st_16x256b_param_5+16];421; CHECK-NEXT:    ld.param.v4.b32 {%r26, %r27, %r28, %r29}, [nvvm_tcgen05_st_16x256b_param_5];422; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x4.b32 [%r1], {%r26, %r27, %r28, %r29, %r22, %r23, %r24, %r25, %r18, %r19, %r20, %r21, %r14, %r15, %r16, %r17};423; CHECK-NEXT:    ld.param.v4.b32 {%r30, %r31, %r32, %r33}, [nvvm_tcgen05_st_16x256b_param_6+112];424; CHECK-NEXT:    ld.param.v4.b32 {%r34, %r35, %r36, %r37}, [nvvm_tcgen05_st_16x256b_param_6+96];425; CHECK-NEXT:    ld.param.v4.b32 {%r38, %r39, %r40, %r41}, [nvvm_tcgen05_st_16x256b_param_6+80];426; CHECK-NEXT:    ld.param.v4.b32 {%r42, %r43, %r44, %r45}, [nvvm_tcgen05_st_16x256b_param_6+64];427; CHECK-NEXT:    ld.param.v4.b32 {%r46, %r47, %r48, %r49}, [nvvm_tcgen05_st_16x256b_param_6+48];428; CHECK-NEXT:    ld.param.v4.b32 {%r50, %r51, %r52, %r53}, [nvvm_tcgen05_st_16x256b_param_6+32];429; CHECK-NEXT:    ld.param.v4.b32 {%r54, %r55, %r56, %r57}, [nvvm_tcgen05_st_16x256b_param_6+16];430; CHECK-NEXT:    ld.param.v4.b32 {%r58, %r59, %r60, %r61}, [nvvm_tcgen05_st_16x256b_param_6];431; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x8.b32 [%r1], {%r58, %r59, %r60, %r61, %r54, %r55, %r56, %r57, %r50, %r51, %r52, %r53, %r46, %r47, %r48, %r49, %r42, %r43, %r44, %r45, %r38, %r39, %r40, %r41, %r34, %r35, %r36, %r37, %r30, %r31, %r32, %r33};432; CHECK-NEXT:    ld.param.v4.b32 {%r62, %r63, %r64, %r65}, [nvvm_tcgen05_st_16x256b_param_7+240];433; CHECK-NEXT:    ld.param.v4.b32 {%r66, %r67, %r68, %r69}, [nvvm_tcgen05_st_16x256b_param_7+224];434; CHECK-NEXT:    ld.param.v4.b32 {%r70, %r71, %r72, %r73}, [nvvm_tcgen05_st_16x256b_param_7+208];435; CHECK-NEXT:    ld.param.v4.b32 {%r74, %r75, %r76, %r77}, [nvvm_tcgen05_st_16x256b_param_7+192];436; CHECK-NEXT:    ld.param.v4.b32 {%r78, %r79, %r80, %r81}, [nvvm_tcgen05_st_16x256b_param_7+176];437; CHECK-NEXT:    ld.param.v4.b32 {%r82, %r83, %r84, %r85}, [nvvm_tcgen05_st_16x256b_param_7+160];438; CHECK-NEXT:    ld.param.v4.b32 {%r86, %r87, %r88, %r89}, [nvvm_tcgen05_st_16x256b_param_7+144];439; CHECK-NEXT:    ld.param.v4.b32 {%r90, %r91, %r92, %r93}, [nvvm_tcgen05_st_16x256b_param_7+128];440; CHECK-NEXT:    ld.param.v4.b32 {%r94, %r95, %r96, %r97}, [nvvm_tcgen05_st_16x256b_param_7+112];441; CHECK-NEXT:    ld.param.v4.b32 {%r98, %r99, %r100, %r101}, [nvvm_tcgen05_st_16x256b_param_7+96];442; CHECK-NEXT:    ld.param.v4.b32 {%r102, %r103, %r104, %r105}, [nvvm_tcgen05_st_16x256b_param_7+80];443; CHECK-NEXT:    ld.param.v4.b32 {%r106, %r107, %r108, %r109}, [nvvm_tcgen05_st_16x256b_param_7+64];444; CHECK-NEXT:    ld.param.v4.b32 {%r110, %r111, %r112, %r113}, [nvvm_tcgen05_st_16x256b_param_7+48];445; CHECK-NEXT:    ld.param.v4.b32 {%r114, %r115, %r116, %r117}, [nvvm_tcgen05_st_16x256b_param_7+32];446; CHECK-NEXT:    ld.param.v4.b32 {%r118, %r119, %r120, %r121}, [nvvm_tcgen05_st_16x256b_param_7+16];447; CHECK-NEXT:    ld.param.v4.b32 {%r122, %r123, %r124, %r125}, [nvvm_tcgen05_st_16x256b_param_7];448; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x16.b32 [%r1], {%r122, %r123, %r124, %r125, %r118, %r119, %r120, %r121, %r114, %r115, %r116, %r117, %r110, %r111, %r112, %r113, %r106, %r107, %r108, %r109, %r102, %r103, %r104, %r105, %r98, %r99, %r100, %r101, %r94, %r95, %r96, %r97, %r90, %r91, %r92, %r93, %r86, %r87, %r88, %r89, %r82, %r83, %r84, %r85, %r78, %r79, %r80, %r81, %r74, %r75, %r76, %r77, %r70, %r71, %r72, %r73, %r66, %r67, %r68, %r69, %r62, %r63, %r64, %r65};449; CHECK-NEXT:    ld.param.v4.b32 {%r126, %r127, %r128, %r129}, [nvvm_tcgen05_st_16x256b_param_8+496];450; CHECK-NEXT:    ld.param.v4.b32 {%r130, %r131, %r132, %r133}, [nvvm_tcgen05_st_16x256b_param_8+480];451; CHECK-NEXT:    ld.param.v4.b32 {%r134, %r135, %r136, %r137}, [nvvm_tcgen05_st_16x256b_param_8+464];452; CHECK-NEXT:    ld.param.v4.b32 {%r138, %r139, %r140, %r141}, [nvvm_tcgen05_st_16x256b_param_8+448];453; CHECK-NEXT:    ld.param.v4.b32 {%r142, %r143, %r144, %r145}, [nvvm_tcgen05_st_16x256b_param_8+432];454; CHECK-NEXT:    ld.param.v4.b32 {%r146, %r147, %r148, %r149}, [nvvm_tcgen05_st_16x256b_param_8+416];455; CHECK-NEXT:    ld.param.v4.b32 {%r150, %r151, %r152, %r153}, [nvvm_tcgen05_st_16x256b_param_8+400];456; CHECK-NEXT:    ld.param.v4.b32 {%r154, %r155, %r156, %r157}, [nvvm_tcgen05_st_16x256b_param_8+384];457; CHECK-NEXT:    ld.param.v4.b32 {%r158, %r159, %r160, %r161}, [nvvm_tcgen05_st_16x256b_param_8+368];458; CHECK-NEXT:    ld.param.v4.b32 {%r162, %r163, %r164, %r165}, [nvvm_tcgen05_st_16x256b_param_8+352];459; CHECK-NEXT:    ld.param.v4.b32 {%r166, %r167, %r168, %r169}, [nvvm_tcgen05_st_16x256b_param_8+336];460; CHECK-NEXT:    ld.param.v4.b32 {%r170, %r171, %r172, %r173}, [nvvm_tcgen05_st_16x256b_param_8+320];461; CHECK-NEXT:    ld.param.v4.b32 {%r174, %r175, %r176, %r177}, [nvvm_tcgen05_st_16x256b_param_8+304];462; CHECK-NEXT:    ld.param.v4.b32 {%r178, %r179, %r180, %r181}, [nvvm_tcgen05_st_16x256b_param_8+288];463; CHECK-NEXT:    ld.param.v4.b32 {%r182, %r183, %r184, %r185}, [nvvm_tcgen05_st_16x256b_param_8+272];464; CHECK-NEXT:    ld.param.v4.b32 {%r186, %r187, %r188, %r189}, [nvvm_tcgen05_st_16x256b_param_8+256];465; CHECK-NEXT:    ld.param.v4.b32 {%r190, %r191, %r192, %r193}, [nvvm_tcgen05_st_16x256b_param_8+240];466; CHECK-NEXT:    ld.param.v4.b32 {%r194, %r195, %r196, %r197}, [nvvm_tcgen05_st_16x256b_param_8+224];467; CHECK-NEXT:    ld.param.v4.b32 {%r198, %r199, %r200, %r201}, [nvvm_tcgen05_st_16x256b_param_8+208];468; CHECK-NEXT:    ld.param.v4.b32 {%r202, %r203, %r204, %r205}, [nvvm_tcgen05_st_16x256b_param_8+192];469; CHECK-NEXT:    ld.param.v4.b32 {%r206, %r207, %r208, %r209}, [nvvm_tcgen05_st_16x256b_param_8+176];470; CHECK-NEXT:    ld.param.v4.b32 {%r210, %r211, %r212, %r213}, [nvvm_tcgen05_st_16x256b_param_8+160];471; CHECK-NEXT:    ld.param.v4.b32 {%r214, %r215, %r216, %r217}, [nvvm_tcgen05_st_16x256b_param_8+144];472; CHECK-NEXT:    ld.param.v4.b32 {%r218, %r219, %r220, %r221}, [nvvm_tcgen05_st_16x256b_param_8+128];473; CHECK-NEXT:    ld.param.v4.b32 {%r222, %r223, %r224, %r225}, [nvvm_tcgen05_st_16x256b_param_8+112];474; CHECK-NEXT:    ld.param.v4.b32 {%r226, %r227, %r228, %r229}, [nvvm_tcgen05_st_16x256b_param_8+96];475; CHECK-NEXT:    ld.param.v4.b32 {%r230, %r231, %r232, %r233}, [nvvm_tcgen05_st_16x256b_param_8+80];476; CHECK-NEXT:    ld.param.v4.b32 {%r234, %r235, %r236, %r237}, [nvvm_tcgen05_st_16x256b_param_8+64];477; CHECK-NEXT:    ld.param.v4.b32 {%r238, %r239, %r240, %r241}, [nvvm_tcgen05_st_16x256b_param_8+48];478; CHECK-NEXT:    ld.param.v4.b32 {%r242, %r243, %r244, %r245}, [nvvm_tcgen05_st_16x256b_param_8+32];479; CHECK-NEXT:    ld.param.v4.b32 {%r246, %r247, %r248, %r249}, [nvvm_tcgen05_st_16x256b_param_8+16];480; CHECK-NEXT:    ld.param.v4.b32 {%r250, %r251, %r252, %r253}, [nvvm_tcgen05_st_16x256b_param_8];481; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x32.b32 [%r1], {%r250, %r251, %r252, %r253, %r246, %r247, %r248, %r249, %r242, %r243, %r244, %r245, %r238, %r239, %r240, %r241, %r234, %r235, %r236, %r237, %r230, %r231, %r232, %r233, %r226, %r227, %r228, %r229, %r222, %r223, %r224, %r225, %r218, %r219, %r220, %r221, %r214, %r215, %r216, %r217, %r210, %r211, %r212, %r213, %r206, %r207, %r208, %r209, %r202, %r203, %r204, %r205, %r198, %r199, %r200, %r201, %r194, %r195, %r196, %r197, %r190, %r191, %r192, %r193, %r186, %r187, %r188, %r189, %r182, %r183, %r184, %r185, %r178, %r179, %r180, %r181, %r174, %r175, %r176, %r177, %r170, %r171, %r172, %r173, %r166, %r167, %r168, %r169, %r162, %r163, %r164, %r165, %r158, %r159, %r160, %r161, %r154, %r155, %r156, %r157, %r150, %r151, %r152, %r153, %r146, %r147, %r148, %r149, %r142, %r143, %r144, %r145, %r138, %r139, %r140, %r141, %r134, %r135, %r136, %r137, %r130, %r131, %r132, %r133, %r126, %r127, %r128, %r129};482; CHECK-NEXT:    ret;483  tail call void @llvm.nvvm.tcgen05.st.16x256b.x1(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 0)484 485  tail call void @llvm.nvvm.tcgen05.st.16x256b.x2(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 0)486 487  tail call void @llvm.nvvm.tcgen05.st.16x256b.x4(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 0)488 489  tail call void @llvm.nvvm.tcgen05.st.16x256b.x8(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 0)490 491  tail call void @llvm.nvvm.tcgen05.st.16x256b.x16(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 0)492 493  tail call void @llvm.nvvm.tcgen05.st.16x256b.x32(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 0)494  ret void495}496 497; CHECK-LABEL: nvvm_tcgen05_st_16x256b_unpack498define void @nvvm_tcgen05_st_16x256b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {499; CHECK-LABEL: nvvm_tcgen05_st_16x256b_unpack(500; CHECK:       {501; CHECK-NEXT:    .reg .b32 %r<254>;502; CHECK-EMPTY:503; CHECK-NEXT:  // %bb.0:504; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x256b_unpack_param_0];505; CHECK-NEXT:    ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [nvvm_tcgen05_st_16x256b_unpack_param_3];506; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x1.unpack::16b.b32 [%r1], {%r2, %r3, %r4, %r5};507; CHECK-NEXT:    ld.param.v4.b32 {%r6, %r7, %r8, %r9}, [nvvm_tcgen05_st_16x256b_unpack_param_4+16];508; CHECK-NEXT:    ld.param.v4.b32 {%r10, %r11, %r12, %r13}, [nvvm_tcgen05_st_16x256b_unpack_param_4];509; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x2.unpack::16b.b32 [%r1], {%r10, %r11, %r12, %r13, %r6, %r7, %r8, %r9};510; CHECK-NEXT:    ld.param.v4.b32 {%r14, %r15, %r16, %r17}, [nvvm_tcgen05_st_16x256b_unpack_param_5+48];511; CHECK-NEXT:    ld.param.v4.b32 {%r18, %r19, %r20, %r21}, [nvvm_tcgen05_st_16x256b_unpack_param_5+32];512; CHECK-NEXT:    ld.param.v4.b32 {%r22, %r23, %r24, %r25}, [nvvm_tcgen05_st_16x256b_unpack_param_5+16];513; CHECK-NEXT:    ld.param.v4.b32 {%r26, %r27, %r28, %r29}, [nvvm_tcgen05_st_16x256b_unpack_param_5];514; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x4.unpack::16b.b32 [%r1], {%r26, %r27, %r28, %r29, %r22, %r23, %r24, %r25, %r18, %r19, %r20, %r21, %r14, %r15, %r16, %r17};515; CHECK-NEXT:    ld.param.v4.b32 {%r30, %r31, %r32, %r33}, [nvvm_tcgen05_st_16x256b_unpack_param_6+112];516; CHECK-NEXT:    ld.param.v4.b32 {%r34, %r35, %r36, %r37}, [nvvm_tcgen05_st_16x256b_unpack_param_6+96];517; CHECK-NEXT:    ld.param.v4.b32 {%r38, %r39, %r40, %r41}, [nvvm_tcgen05_st_16x256b_unpack_param_6+80];518; CHECK-NEXT:    ld.param.v4.b32 {%r42, %r43, %r44, %r45}, [nvvm_tcgen05_st_16x256b_unpack_param_6+64];519; CHECK-NEXT:    ld.param.v4.b32 {%r46, %r47, %r48, %r49}, [nvvm_tcgen05_st_16x256b_unpack_param_6+48];520; CHECK-NEXT:    ld.param.v4.b32 {%r50, %r51, %r52, %r53}, [nvvm_tcgen05_st_16x256b_unpack_param_6+32];521; CHECK-NEXT:    ld.param.v4.b32 {%r54, %r55, %r56, %r57}, [nvvm_tcgen05_st_16x256b_unpack_param_6+16];522; CHECK-NEXT:    ld.param.v4.b32 {%r58, %r59, %r60, %r61}, [nvvm_tcgen05_st_16x256b_unpack_param_6];523; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x8.unpack::16b.b32 [%r1], {%r58, %r59, %r60, %r61, %r54, %r55, %r56, %r57, %r50, %r51, %r52, %r53, %r46, %r47, %r48, %r49, %r42, %r43, %r44, %r45, %r38, %r39, %r40, %r41, %r34, %r35, %r36, %r37, %r30, %r31, %r32, %r33};524; CHECK-NEXT:    ld.param.v4.b32 {%r62, %r63, %r64, %r65}, [nvvm_tcgen05_st_16x256b_unpack_param_7+240];525; CHECK-NEXT:    ld.param.v4.b32 {%r66, %r67, %r68, %r69}, [nvvm_tcgen05_st_16x256b_unpack_param_7+224];526; CHECK-NEXT:    ld.param.v4.b32 {%r70, %r71, %r72, %r73}, [nvvm_tcgen05_st_16x256b_unpack_param_7+208];527; CHECK-NEXT:    ld.param.v4.b32 {%r74, %r75, %r76, %r77}, [nvvm_tcgen05_st_16x256b_unpack_param_7+192];528; CHECK-NEXT:    ld.param.v4.b32 {%r78, %r79, %r80, %r81}, [nvvm_tcgen05_st_16x256b_unpack_param_7+176];529; CHECK-NEXT:    ld.param.v4.b32 {%r82, %r83, %r84, %r85}, [nvvm_tcgen05_st_16x256b_unpack_param_7+160];530; CHECK-NEXT:    ld.param.v4.b32 {%r86, %r87, %r88, %r89}, [nvvm_tcgen05_st_16x256b_unpack_param_7+144];531; CHECK-NEXT:    ld.param.v4.b32 {%r90, %r91, %r92, %r93}, [nvvm_tcgen05_st_16x256b_unpack_param_7+128];532; CHECK-NEXT:    ld.param.v4.b32 {%r94, %r95, %r96, %r97}, [nvvm_tcgen05_st_16x256b_unpack_param_7+112];533; CHECK-NEXT:    ld.param.v4.b32 {%r98, %r99, %r100, %r101}, [nvvm_tcgen05_st_16x256b_unpack_param_7+96];534; CHECK-NEXT:    ld.param.v4.b32 {%r102, %r103, %r104, %r105}, [nvvm_tcgen05_st_16x256b_unpack_param_7+80];535; CHECK-NEXT:    ld.param.v4.b32 {%r106, %r107, %r108, %r109}, [nvvm_tcgen05_st_16x256b_unpack_param_7+64];536; CHECK-NEXT:    ld.param.v4.b32 {%r110, %r111, %r112, %r113}, [nvvm_tcgen05_st_16x256b_unpack_param_7+48];537; CHECK-NEXT:    ld.param.v4.b32 {%r114, %r115, %r116, %r117}, [nvvm_tcgen05_st_16x256b_unpack_param_7+32];538; CHECK-NEXT:    ld.param.v4.b32 {%r118, %r119, %r120, %r121}, [nvvm_tcgen05_st_16x256b_unpack_param_7+16];539; CHECK-NEXT:    ld.param.v4.b32 {%r122, %r123, %r124, %r125}, [nvvm_tcgen05_st_16x256b_unpack_param_7];540; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x16.unpack::16b.b32 [%r1], {%r122, %r123, %r124, %r125, %r118, %r119, %r120, %r121, %r114, %r115, %r116, %r117, %r110, %r111, %r112, %r113, %r106, %r107, %r108, %r109, %r102, %r103, %r104, %r105, %r98, %r99, %r100, %r101, %r94, %r95, %r96, %r97, %r90, %r91, %r92, %r93, %r86, %r87, %r88, %r89, %r82, %r83, %r84, %r85, %r78, %r79, %r80, %r81, %r74, %r75, %r76, %r77, %r70, %r71, %r72, %r73, %r66, %r67, %r68, %r69, %r62, %r63, %r64, %r65};541; CHECK-NEXT:    ld.param.v4.b32 {%r126, %r127, %r128, %r129}, [nvvm_tcgen05_st_16x256b_unpack_param_8+496];542; CHECK-NEXT:    ld.param.v4.b32 {%r130, %r131, %r132, %r133}, [nvvm_tcgen05_st_16x256b_unpack_param_8+480];543; CHECK-NEXT:    ld.param.v4.b32 {%r134, %r135, %r136, %r137}, [nvvm_tcgen05_st_16x256b_unpack_param_8+464];544; CHECK-NEXT:    ld.param.v4.b32 {%r138, %r139, %r140, %r141}, [nvvm_tcgen05_st_16x256b_unpack_param_8+448];545; CHECK-NEXT:    ld.param.v4.b32 {%r142, %r143, %r144, %r145}, [nvvm_tcgen05_st_16x256b_unpack_param_8+432];546; CHECK-NEXT:    ld.param.v4.b32 {%r146, %r147, %r148, %r149}, [nvvm_tcgen05_st_16x256b_unpack_param_8+416];547; CHECK-NEXT:    ld.param.v4.b32 {%r150, %r151, %r152, %r153}, [nvvm_tcgen05_st_16x256b_unpack_param_8+400];548; CHECK-NEXT:    ld.param.v4.b32 {%r154, %r155, %r156, %r157}, [nvvm_tcgen05_st_16x256b_unpack_param_8+384];549; CHECK-NEXT:    ld.param.v4.b32 {%r158, %r159, %r160, %r161}, [nvvm_tcgen05_st_16x256b_unpack_param_8+368];550; CHECK-NEXT:    ld.param.v4.b32 {%r162, %r163, %r164, %r165}, [nvvm_tcgen05_st_16x256b_unpack_param_8+352];551; CHECK-NEXT:    ld.param.v4.b32 {%r166, %r167, %r168, %r169}, [nvvm_tcgen05_st_16x256b_unpack_param_8+336];552; CHECK-NEXT:    ld.param.v4.b32 {%r170, %r171, %r172, %r173}, [nvvm_tcgen05_st_16x256b_unpack_param_8+320];553; CHECK-NEXT:    ld.param.v4.b32 {%r174, %r175, %r176, %r177}, [nvvm_tcgen05_st_16x256b_unpack_param_8+304];554; CHECK-NEXT:    ld.param.v4.b32 {%r178, %r179, %r180, %r181}, [nvvm_tcgen05_st_16x256b_unpack_param_8+288];555; CHECK-NEXT:    ld.param.v4.b32 {%r182, %r183, %r184, %r185}, [nvvm_tcgen05_st_16x256b_unpack_param_8+272];556; CHECK-NEXT:    ld.param.v4.b32 {%r186, %r187, %r188, %r189}, [nvvm_tcgen05_st_16x256b_unpack_param_8+256];557; CHECK-NEXT:    ld.param.v4.b32 {%r190, %r191, %r192, %r193}, [nvvm_tcgen05_st_16x256b_unpack_param_8+240];558; CHECK-NEXT:    ld.param.v4.b32 {%r194, %r195, %r196, %r197}, [nvvm_tcgen05_st_16x256b_unpack_param_8+224];559; CHECK-NEXT:    ld.param.v4.b32 {%r198, %r199, %r200, %r201}, [nvvm_tcgen05_st_16x256b_unpack_param_8+208];560; CHECK-NEXT:    ld.param.v4.b32 {%r202, %r203, %r204, %r205}, [nvvm_tcgen05_st_16x256b_unpack_param_8+192];561; CHECK-NEXT:    ld.param.v4.b32 {%r206, %r207, %r208, %r209}, [nvvm_tcgen05_st_16x256b_unpack_param_8+176];562; CHECK-NEXT:    ld.param.v4.b32 {%r210, %r211, %r212, %r213}, [nvvm_tcgen05_st_16x256b_unpack_param_8+160];563; CHECK-NEXT:    ld.param.v4.b32 {%r214, %r215, %r216, %r217}, [nvvm_tcgen05_st_16x256b_unpack_param_8+144];564; CHECK-NEXT:    ld.param.v4.b32 {%r218, %r219, %r220, %r221}, [nvvm_tcgen05_st_16x256b_unpack_param_8+128];565; CHECK-NEXT:    ld.param.v4.b32 {%r222, %r223, %r224, %r225}, [nvvm_tcgen05_st_16x256b_unpack_param_8+112];566; CHECK-NEXT:    ld.param.v4.b32 {%r226, %r227, %r228, %r229}, [nvvm_tcgen05_st_16x256b_unpack_param_8+96];567; CHECK-NEXT:    ld.param.v4.b32 {%r230, %r231, %r232, %r233}, [nvvm_tcgen05_st_16x256b_unpack_param_8+80];568; CHECK-NEXT:    ld.param.v4.b32 {%r234, %r235, %r236, %r237}, [nvvm_tcgen05_st_16x256b_unpack_param_8+64];569; CHECK-NEXT:    ld.param.v4.b32 {%r238, %r239, %r240, %r241}, [nvvm_tcgen05_st_16x256b_unpack_param_8+48];570; CHECK-NEXT:    ld.param.v4.b32 {%r242, %r243, %r244, %r245}, [nvvm_tcgen05_st_16x256b_unpack_param_8+32];571; CHECK-NEXT:    ld.param.v4.b32 {%r246, %r247, %r248, %r249}, [nvvm_tcgen05_st_16x256b_unpack_param_8+16];572; CHECK-NEXT:    ld.param.v4.b32 {%r250, %r251, %r252, %r253}, [nvvm_tcgen05_st_16x256b_unpack_param_8];573; CHECK-NEXT:    tcgen05.st.sync.aligned.16x256b.x32.unpack::16b.b32 [%r1], {%r250, %r251, %r252, %r253, %r246, %r247, %r248, %r249, %r242, %r243, %r244, %r245, %r238, %r239, %r240, %r241, %r234, %r235, %r236, %r237, %r230, %r231, %r232, %r233, %r226, %r227, %r228, %r229, %r222, %r223, %r224, %r225, %r218, %r219, %r220, %r221, %r214, %r215, %r216, %r217, %r210, %r211, %r212, %r213, %r206, %r207, %r208, %r209, %r202, %r203, %r204, %r205, %r198, %r199, %r200, %r201, %r194, %r195, %r196, %r197, %r190, %r191, %r192, %r193, %r186, %r187, %r188, %r189, %r182, %r183, %r184, %r185, %r178, %r179, %r180, %r181, %r174, %r175, %r176, %r177, %r170, %r171, %r172, %r173, %r166, %r167, %r168, %r169, %r162, %r163, %r164, %r165, %r158, %r159, %r160, %r161, %r154, %r155, %r156, %r157, %r150, %r151, %r152, %r153, %r146, %r147, %r148, %r149, %r142, %r143, %r144, %r145, %r138, %r139, %r140, %r141, %r134, %r135, %r136, %r137, %r130, %r131, %r132, %r133, %r126, %r127, %r128, %r129};574; CHECK-NEXT:    ret;575  tail call void @llvm.nvvm.tcgen05.st.16x256b.x1(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 1)576 577  tail call void @llvm.nvvm.tcgen05.st.16x256b.x2(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 1)578 579  tail call void @llvm.nvvm.tcgen05.st.16x256b.x4(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 1)580 581  tail call void @llvm.nvvm.tcgen05.st.16x256b.x8(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 1)582 583  tail call void @llvm.nvvm.tcgen05.st.16x256b.x16(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 1)584 585  tail call void @llvm.nvvm.tcgen05.st.16x256b.x32(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 1)586  ret void587}588 589; CHECK-LABEL: nvvm_tcgen05_st_32x32b590define void @nvvm_tcgen05_st_32x32b(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {591; CHECK-LABEL: nvvm_tcgen05_st_32x32b(592; CHECK:       {593; CHECK-NEXT:    .reg .b32 %r<257>;594; CHECK-EMPTY:595; CHECK-NEXT:  // %bb.0:596; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_32x32b_param_0];597; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_32x32b_param_1];598; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x1.b32 [%r1], {%r2};599; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_32x32b_param_2];600; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x2.b32 [%r1], {%r3, %r4};601; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_32x32b_param_3];602; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x4.b32 [%r1], {%r5, %r6, %r7, %r8};603; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_32x32b_param_4+16];604; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_32x32b_param_4];605; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x8.b32 [%r1], {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};606; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_32x32b_param_5+48];607; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_32x32b_param_5+32];608; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_32x32b_param_5+16];609; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_32x32b_param_5];610; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x16.b32 [%r1], {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};611; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_32x32b_param_6+112];612; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_32x32b_param_6+96];613; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_32x32b_param_6+80];614; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_32x32b_param_6+64];615; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_32x32b_param_6+48];616; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_32x32b_param_6+32];617; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_32x32b_param_6+16];618; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_32x32b_param_6];619; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x32.b32 [%r1], {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};620; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_32x32b_param_7+240];621; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_32x32b_param_7+224];622; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_32x32b_param_7+208];623; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_32x32b_param_7+192];624; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_32x32b_param_7+176];625; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_32x32b_param_7+160];626; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_32x32b_param_7+144];627; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_32x32b_param_7+128];628; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_32x32b_param_7+112];629; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_32x32b_param_7+96];630; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_32x32b_param_7+80];631; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_32x32b_param_7+64];632; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_32x32b_param_7+48];633; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_32x32b_param_7+32];634; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_32x32b_param_7+16];635; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_32x32b_param_7];636; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x64.b32 [%r1], {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};637; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_32x32b_param_8+496];638; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_32x32b_param_8+480];639; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_32x32b_param_8+464];640; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_32x32b_param_8+448];641; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_32x32b_param_8+432];642; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_32x32b_param_8+416];643; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_32x32b_param_8+400];644; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_32x32b_param_8+384];645; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_32x32b_param_8+368];646; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_32x32b_param_8+352];647; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_32x32b_param_8+336];648; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_32x32b_param_8+320];649; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_32x32b_param_8+304];650; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_32x32b_param_8+288];651; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_32x32b_param_8+272];652; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_32x32b_param_8+256];653; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_32x32b_param_8+240];654; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_32x32b_param_8+224];655; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_32x32b_param_8+208];656; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_32x32b_param_8+192];657; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_32x32b_param_8+176];658; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_32x32b_param_8+160];659; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_32x32b_param_8+144];660; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_32x32b_param_8+128];661; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_32x32b_param_8+112];662; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_32x32b_param_8+96];663; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_32x32b_param_8+80];664; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_32x32b_param_8+64];665; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_32x32b_param_8+48];666; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_32x32b_param_8+32];667; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_32x32b_param_8+16];668; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_32x32b_param_8];669; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x128.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};670; CHECK-NEXT:    ret;671  tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 0)672 673  tail call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 0)674 675  tail call void @llvm.nvvm.tcgen05.st.32x32b.x4(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 0)676 677  tail call void @llvm.nvvm.tcgen05.st.32x32b.x8(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 0)678 679  tail call void @llvm.nvvm.tcgen05.st.32x32b.x16(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 0)680 681  tail call void @llvm.nvvm.tcgen05.st.32x32b.x32(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 0)682 683  tail call void @llvm.nvvm.tcgen05.st.32x32b.x64(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 0)684 685  tail call void @llvm.nvvm.tcgen05.st.32x32b.x128(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 0)686  ret void687}688 689; CHECK-LABEL: nvvm_tcgen05_st_32x32b_unpack690define void @nvvm_tcgen05_st_32x32b_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {691; CHECK-LABEL: nvvm_tcgen05_st_32x32b_unpack(692; CHECK:       {693; CHECK-NEXT:    .reg .b32 %r<257>;694; CHECK-EMPTY:695; CHECK-NEXT:  // %bb.0:696; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_32x32b_unpack_param_0];697; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_32x32b_unpack_param_1];698; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x1.unpack::16b.b32 [%r1], {%r2};699; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_32x32b_unpack_param_2];700; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x2.unpack::16b.b32 [%r1], {%r3, %r4};701; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_32x32b_unpack_param_3];702; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x4.unpack::16b.b32 [%r1], {%r5, %r6, %r7, %r8};703; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_32x32b_unpack_param_4+16];704; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_32x32b_unpack_param_4];705; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x8.unpack::16b.b32 [%r1], {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};706; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_32x32b_unpack_param_5+48];707; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_32x32b_unpack_param_5+32];708; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_32x32b_unpack_param_5+16];709; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_32x32b_unpack_param_5];710; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x16.unpack::16b.b32 [%r1], {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};711; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_32x32b_unpack_param_6+112];712; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_32x32b_unpack_param_6+96];713; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_32x32b_unpack_param_6+80];714; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_32x32b_unpack_param_6+64];715; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_32x32b_unpack_param_6+48];716; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_32x32b_unpack_param_6+32];717; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_32x32b_unpack_param_6+16];718; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_32x32b_unpack_param_6];719; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x32.unpack::16b.b32 [%r1], {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};720; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_32x32b_unpack_param_7+240];721; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_32x32b_unpack_param_7+224];722; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_32x32b_unpack_param_7+208];723; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_32x32b_unpack_param_7+192];724; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_32x32b_unpack_param_7+176];725; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_32x32b_unpack_param_7+160];726; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_32x32b_unpack_param_7+144];727; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_32x32b_unpack_param_7+128];728; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_32x32b_unpack_param_7+112];729; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_32x32b_unpack_param_7+96];730; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_32x32b_unpack_param_7+80];731; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_32x32b_unpack_param_7+64];732; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_32x32b_unpack_param_7+48];733; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_32x32b_unpack_param_7+32];734; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_32x32b_unpack_param_7+16];735; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_32x32b_unpack_param_7];736; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x64.unpack::16b.b32 [%r1], {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};737; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_32x32b_unpack_param_8+496];738; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_32x32b_unpack_param_8+480];739; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_32x32b_unpack_param_8+464];740; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_32x32b_unpack_param_8+448];741; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_32x32b_unpack_param_8+432];742; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_32x32b_unpack_param_8+416];743; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_32x32b_unpack_param_8+400];744; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_32x32b_unpack_param_8+384];745; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_32x32b_unpack_param_8+368];746; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_32x32b_unpack_param_8+352];747; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_32x32b_unpack_param_8+336];748; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_32x32b_unpack_param_8+320];749; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_32x32b_unpack_param_8+304];750; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_32x32b_unpack_param_8+288];751; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_32x32b_unpack_param_8+272];752; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_32x32b_unpack_param_8+256];753; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_32x32b_unpack_param_8+240];754; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_32x32b_unpack_param_8+224];755; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_32x32b_unpack_param_8+208];756; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_32x32b_unpack_param_8+192];757; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_32x32b_unpack_param_8+176];758; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_32x32b_unpack_param_8+160];759; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_32x32b_unpack_param_8+144];760; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_32x32b_unpack_param_8+128];761; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_32x32b_unpack_param_8+112];762; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_32x32b_unpack_param_8+96];763; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_32x32b_unpack_param_8+80];764; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_32x32b_unpack_param_8+64];765; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_32x32b_unpack_param_8+48];766; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_32x32b_unpack_param_8+32];767; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_32x32b_unpack_param_8+16];768; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_32x32b_unpack_param_8];769; CHECK-NEXT:    tcgen05.st.sync.aligned.32x32b.x128.unpack::16b.b32 [%r1], {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};770; CHECK-NEXT:    ret;771  tail call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) %taddr, i32 %stv1, i1 1)772 773  tail call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) %taddr, <2 x i32> %stv2, i1 1)774 775  tail call void @llvm.nvvm.tcgen05.st.32x32b.x4(ptr addrspace(6) %taddr, <4 x i32> %stv4, i1 1)776 777  tail call void @llvm.nvvm.tcgen05.st.32x32b.x8(ptr addrspace(6) %taddr, <8 x i32> %stv8, i1 1)778 779  tail call void @llvm.nvvm.tcgen05.st.32x32b.x16(ptr addrspace(6) %taddr, <16 x i32> %stv16, i1 1)780 781  tail call void @llvm.nvvm.tcgen05.st.32x32b.x32(ptr addrspace(6) %taddr, <32 x i32> %stv32, i1 1)782 783  tail call void @llvm.nvvm.tcgen05.st.32x32b.x64(ptr addrspace(6) %taddr, <64 x i32> %stv64, i1 1)784 785  tail call void @llvm.nvvm.tcgen05.st.32x32b.x128(ptr addrspace(6) %taddr, <128 x i32> %stv128, i1 1)786  ret void787}788 789; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2790define void @nvvm_tcgen05_st_16x32bx2(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {791; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2(792; CHECK:       {793; CHECK-NEXT:    .reg .b32 %r<257>;794; CHECK-EMPTY:795; CHECK-NEXT:  // %bb.0:796; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x32bx2_param_0];797; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_16x32bx2_param_1];798; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x1.b32 [%r1], 2, {%r2};799; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_16x32bx2_param_2];800; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x2.b32 [%r1], 2, {%r3, %r4};801; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_16x32bx2_param_3];802; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x4.b32 [%r1], 2, {%r5, %r6, %r7, %r8};803; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_16x32bx2_param_4+16];804; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_16x32bx2_param_4];805; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x8.b32 [%r1], 2, {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};806; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_16x32bx2_param_5+48];807; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_16x32bx2_param_5+32];808; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_16x32bx2_param_5+16];809; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_16x32bx2_param_5];810; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x16.b32 [%r1], 2, {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};811; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_16x32bx2_param_6+112];812; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_16x32bx2_param_6+96];813; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_16x32bx2_param_6+80];814; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_16x32bx2_param_6+64];815; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_16x32bx2_param_6+48];816; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_16x32bx2_param_6+32];817; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_16x32bx2_param_6+16];818; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_16x32bx2_param_6];819; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x32.b32 [%r1], 2, {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};820; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_16x32bx2_param_7+240];821; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_16x32bx2_param_7+224];822; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_16x32bx2_param_7+208];823; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_16x32bx2_param_7+192];824; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_16x32bx2_param_7+176];825; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_16x32bx2_param_7+160];826; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_16x32bx2_param_7+144];827; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_16x32bx2_param_7+128];828; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_16x32bx2_param_7+112];829; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_16x32bx2_param_7+96];830; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_16x32bx2_param_7+80];831; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_16x32bx2_param_7+64];832; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_16x32bx2_param_7+48];833; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_16x32bx2_param_7+32];834; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_16x32bx2_param_7+16];835; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_16x32bx2_param_7];836; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x64.b32 [%r1], 2, {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};837; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_16x32bx2_param_8+496];838; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_16x32bx2_param_8+480];839; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_16x32bx2_param_8+464];840; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_16x32bx2_param_8+448];841; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_16x32bx2_param_8+432];842; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_16x32bx2_param_8+416];843; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_16x32bx2_param_8+400];844; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_16x32bx2_param_8+384];845; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_16x32bx2_param_8+368];846; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_16x32bx2_param_8+352];847; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_16x32bx2_param_8+336];848; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_16x32bx2_param_8+320];849; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_16x32bx2_param_8+304];850; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_16x32bx2_param_8+288];851; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_16x32bx2_param_8+272];852; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_16x32bx2_param_8+256];853; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_16x32bx2_param_8+240];854; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_16x32bx2_param_8+224];855; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_16x32bx2_param_8+208];856; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_16x32bx2_param_8+192];857; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_16x32bx2_param_8+176];858; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_16x32bx2_param_8+160];859; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_16x32bx2_param_8+144];860; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_16x32bx2_param_8+128];861; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_16x32bx2_param_8+112];862; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_16x32bx2_param_8+96];863; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_16x32bx2_param_8+80];864; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_16x32bx2_param_8+64];865; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_16x32bx2_param_8+48];866; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_16x32bx2_param_8+32];867; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_16x32bx2_param_8+16];868; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x32bx2_param_8];869; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x128.b32 [%r1], 2, {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};870; CHECK-NEXT:    ret;871  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i32 %stv1, i1 0)872 873  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, <2 x i32> %stv2, i1 0)874 875  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x4(ptr addrspace(6) %taddr, i64 2, <4 x i32> %stv4, i1 0)876 877  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x8(ptr addrspace(6) %taddr, i64 2, <8 x i32> %stv8, i1 0)878 879  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x16(ptr addrspace(6) %taddr, i64 2, <16 x i32> %stv16, i1 0)880 881  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x32(ptr addrspace(6) %taddr, i64 2, <32 x i32> %stv32, i1 0)882 883  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x64(ptr addrspace(6) %taddr, i64 2, <64 x i32> %stv64, i1 0)884 885  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x128(ptr addrspace(6) %taddr, i64 2, <128 x i32> %stv128, i1 0)886  ret void887}888 889; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2_unpack890define void @nvvm_tcgen05_st_16x32bx2_unpack(ptr addrspace(6) %taddr, i32 %stv1, <2 x i32> %stv2, <4 x i32> %stv4, <8 x i32> %stv8, <16 x i32> %stv16, <32 x i32> %stv32, <64 x i32> %stv64, <128 x i32> %stv128) {891; CHECK-LABEL: nvvm_tcgen05_st_16x32bx2_unpack(892; CHECK:       {893; CHECK-NEXT:    .reg .b32 %r<257>;894; CHECK-EMPTY:895; CHECK-NEXT:  // %bb.0:896; CHECK-NEXT:    ld.param.b32 %r1, [nvvm_tcgen05_st_16x32bx2_unpack_param_0];897; CHECK-NEXT:    ld.param.b32 %r2, [nvvm_tcgen05_st_16x32bx2_unpack_param_1];898; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x1.unpack::16b.b32 [%r1], 2, {%r2};899; CHECK-NEXT:    ld.param.v2.b32 {%r3, %r4}, [nvvm_tcgen05_st_16x32bx2_unpack_param_2];900; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x2.unpack::16b.b32 [%r1], 2, {%r3, %r4};901; CHECK-NEXT:    ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [nvvm_tcgen05_st_16x32bx2_unpack_param_3];902; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x4.unpack::16b.b32 [%r1], 2, {%r5, %r6, %r7, %r8};903; CHECK-NEXT:    ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [nvvm_tcgen05_st_16x32bx2_unpack_param_4+16];904; CHECK-NEXT:    ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [nvvm_tcgen05_st_16x32bx2_unpack_param_4];905; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x8.unpack::16b.b32 [%r1], 2, {%r13, %r14, %r15, %r16, %r9, %r10, %r11, %r12};906; CHECK-NEXT:    ld.param.v4.b32 {%r17, %r18, %r19, %r20}, [nvvm_tcgen05_st_16x32bx2_unpack_param_5+48];907; CHECK-NEXT:    ld.param.v4.b32 {%r21, %r22, %r23, %r24}, [nvvm_tcgen05_st_16x32bx2_unpack_param_5+32];908; CHECK-NEXT:    ld.param.v4.b32 {%r25, %r26, %r27, %r28}, [nvvm_tcgen05_st_16x32bx2_unpack_param_5+16];909; CHECK-NEXT:    ld.param.v4.b32 {%r29, %r30, %r31, %r32}, [nvvm_tcgen05_st_16x32bx2_unpack_param_5];910; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x16.unpack::16b.b32 [%r1], 2, {%r29, %r30, %r31, %r32, %r25, %r26, %r27, %r28, %r21, %r22, %r23, %r24, %r17, %r18, %r19, %r20};911; CHECK-NEXT:    ld.param.v4.b32 {%r33, %r34, %r35, %r36}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+112];912; CHECK-NEXT:    ld.param.v4.b32 {%r37, %r38, %r39, %r40}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+96];913; CHECK-NEXT:    ld.param.v4.b32 {%r41, %r42, %r43, %r44}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+80];914; CHECK-NEXT:    ld.param.v4.b32 {%r45, %r46, %r47, %r48}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+64];915; CHECK-NEXT:    ld.param.v4.b32 {%r49, %r50, %r51, %r52}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+48];916; CHECK-NEXT:    ld.param.v4.b32 {%r53, %r54, %r55, %r56}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+32];917; CHECK-NEXT:    ld.param.v4.b32 {%r57, %r58, %r59, %r60}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6+16];918; CHECK-NEXT:    ld.param.v4.b32 {%r61, %r62, %r63, %r64}, [nvvm_tcgen05_st_16x32bx2_unpack_param_6];919; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x32.unpack::16b.b32 [%r1], 2, {%r61, %r62, %r63, %r64, %r57, %r58, %r59, %r60, %r53, %r54, %r55, %r56, %r49, %r50, %r51, %r52, %r45, %r46, %r47, %r48, %r41, %r42, %r43, %r44, %r37, %r38, %r39, %r40, %r33, %r34, %r35, %r36};920; CHECK-NEXT:    ld.param.v4.b32 {%r65, %r66, %r67, %r68}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+240];921; CHECK-NEXT:    ld.param.v4.b32 {%r69, %r70, %r71, %r72}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+224];922; CHECK-NEXT:    ld.param.v4.b32 {%r73, %r74, %r75, %r76}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+208];923; CHECK-NEXT:    ld.param.v4.b32 {%r77, %r78, %r79, %r80}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+192];924; CHECK-NEXT:    ld.param.v4.b32 {%r81, %r82, %r83, %r84}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+176];925; CHECK-NEXT:    ld.param.v4.b32 {%r85, %r86, %r87, %r88}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+160];926; CHECK-NEXT:    ld.param.v4.b32 {%r89, %r90, %r91, %r92}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+144];927; CHECK-NEXT:    ld.param.v4.b32 {%r93, %r94, %r95, %r96}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+128];928; CHECK-NEXT:    ld.param.v4.b32 {%r97, %r98, %r99, %r100}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+112];929; CHECK-NEXT:    ld.param.v4.b32 {%r101, %r102, %r103, %r104}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+96];930; CHECK-NEXT:    ld.param.v4.b32 {%r105, %r106, %r107, %r108}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+80];931; CHECK-NEXT:    ld.param.v4.b32 {%r109, %r110, %r111, %r112}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+64];932; CHECK-NEXT:    ld.param.v4.b32 {%r113, %r114, %r115, %r116}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+48];933; CHECK-NEXT:    ld.param.v4.b32 {%r117, %r118, %r119, %r120}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+32];934; CHECK-NEXT:    ld.param.v4.b32 {%r121, %r122, %r123, %r124}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7+16];935; CHECK-NEXT:    ld.param.v4.b32 {%r125, %r126, %r127, %r128}, [nvvm_tcgen05_st_16x32bx2_unpack_param_7];936; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x64.unpack::16b.b32 [%r1], 2, {%r125, %r126, %r127, %r128, %r121, %r122, %r123, %r124, %r117, %r118, %r119, %r120, %r113, %r114, %r115, %r116, %r109, %r110, %r111, %r112, %r105, %r106, %r107, %r108, %r101, %r102, %r103, %r104, %r97, %r98, %r99, %r100, %r93, %r94, %r95, %r96, %r89, %r90, %r91, %r92, %r85, %r86, %r87, %r88, %r81, %r82, %r83, %r84, %r77, %r78, %r79, %r80, %r73, %r74, %r75, %r76, %r69, %r70, %r71, %r72, %r65, %r66, %r67, %r68};937; CHECK-NEXT:    ld.param.v4.b32 {%r129, %r130, %r131, %r132}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+496];938; CHECK-NEXT:    ld.param.v4.b32 {%r133, %r134, %r135, %r136}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+480];939; CHECK-NEXT:    ld.param.v4.b32 {%r137, %r138, %r139, %r140}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+464];940; CHECK-NEXT:    ld.param.v4.b32 {%r141, %r142, %r143, %r144}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+448];941; CHECK-NEXT:    ld.param.v4.b32 {%r145, %r146, %r147, %r148}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+432];942; CHECK-NEXT:    ld.param.v4.b32 {%r149, %r150, %r151, %r152}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+416];943; CHECK-NEXT:    ld.param.v4.b32 {%r153, %r154, %r155, %r156}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+400];944; CHECK-NEXT:    ld.param.v4.b32 {%r157, %r158, %r159, %r160}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+384];945; CHECK-NEXT:    ld.param.v4.b32 {%r161, %r162, %r163, %r164}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+368];946; CHECK-NEXT:    ld.param.v4.b32 {%r165, %r166, %r167, %r168}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+352];947; CHECK-NEXT:    ld.param.v4.b32 {%r169, %r170, %r171, %r172}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+336];948; CHECK-NEXT:    ld.param.v4.b32 {%r173, %r174, %r175, %r176}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+320];949; CHECK-NEXT:    ld.param.v4.b32 {%r177, %r178, %r179, %r180}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+304];950; CHECK-NEXT:    ld.param.v4.b32 {%r181, %r182, %r183, %r184}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+288];951; CHECK-NEXT:    ld.param.v4.b32 {%r185, %r186, %r187, %r188}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+272];952; CHECK-NEXT:    ld.param.v4.b32 {%r189, %r190, %r191, %r192}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+256];953; CHECK-NEXT:    ld.param.v4.b32 {%r193, %r194, %r195, %r196}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+240];954; CHECK-NEXT:    ld.param.v4.b32 {%r197, %r198, %r199, %r200}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+224];955; CHECK-NEXT:    ld.param.v4.b32 {%r201, %r202, %r203, %r204}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+208];956; CHECK-NEXT:    ld.param.v4.b32 {%r205, %r206, %r207, %r208}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+192];957; CHECK-NEXT:    ld.param.v4.b32 {%r209, %r210, %r211, %r212}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+176];958; CHECK-NEXT:    ld.param.v4.b32 {%r213, %r214, %r215, %r216}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+160];959; CHECK-NEXT:    ld.param.v4.b32 {%r217, %r218, %r219, %r220}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+144];960; CHECK-NEXT:    ld.param.v4.b32 {%r221, %r222, %r223, %r224}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+128];961; CHECK-NEXT:    ld.param.v4.b32 {%r225, %r226, %r227, %r228}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+112];962; CHECK-NEXT:    ld.param.v4.b32 {%r229, %r230, %r231, %r232}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+96];963; CHECK-NEXT:    ld.param.v4.b32 {%r233, %r234, %r235, %r236}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+80];964; CHECK-NEXT:    ld.param.v4.b32 {%r237, %r238, %r239, %r240}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+64];965; CHECK-NEXT:    ld.param.v4.b32 {%r241, %r242, %r243, %r244}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+48];966; CHECK-NEXT:    ld.param.v4.b32 {%r245, %r246, %r247, %r248}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+32];967; CHECK-NEXT:    ld.param.v4.b32 {%r249, %r250, %r251, %r252}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8+16];968; CHECK-NEXT:    ld.param.v4.b32 {%r253, %r254, %r255, %r256}, [nvvm_tcgen05_st_16x32bx2_unpack_param_8];969; CHECK-NEXT:    tcgen05.st.sync.aligned.16x32bx2.x128.unpack::16b.b32 [%r1], 2, {%r253, %r254, %r255, %r256, %r249, %r250, %r251, %r252, %r245, %r246, %r247, %r248, %r241, %r242, %r243, %r244, %r237, %r238, %r239, %r240, %r233, %r234, %r235, %r236, %r229, %r230, %r231, %r232, %r225, %r226, %r227, %r228, %r221, %r222, %r223, %r224, %r217, %r218, %r219, %r220, %r213, %r214, %r215, %r216, %r209, %r210, %r211, %r212, %r205, %r206, %r207, %r208, %r201, %r202, %r203, %r204, %r197, %r198, %r199, %r200, %r193, %r194, %r195, %r196, %r189, %r190, %r191, %r192, %r185, %r186, %r187, %r188, %r181, %r182, %r183, %r184, %r177, %r178, %r179, %r180, %r173, %r174, %r175, %r176, %r169, %r170, %r171, %r172, %r165, %r166, %r167, %r168, %r161, %r162, %r163, %r164, %r157, %r158, %r159, %r160, %r153, %r154, %r155, %r156, %r149, %r150, %r151, %r152, %r145, %r146, %r147, %r148, %r141, %r142, %r143, %r144, %r137, %r138, %r139, %r140, %r133, %r134, %r135, %r136, %r129, %r130, %r131, %r132};970; CHECK-NEXT:    ret;971  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) %taddr, i64 2, i32 %stv1, i1 1)972 973  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) %taddr, i64 2, <2 x i32> %stv2, i1 1)974 975  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x4(ptr addrspace(6) %taddr, i64 2, <4 x i32> %stv4, i1 1)976 977  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x8(ptr addrspace(6) %taddr, i64 2, <8 x i32> %stv8, i1 1)978 979  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x16(ptr addrspace(6) %taddr, i64 2, <16 x i32> %stv16, i1 1)980 981  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x32(ptr addrspace(6) %taddr, i64 2, <32 x i32> %stv32, i1 1)982 983  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x64(ptr addrspace(6) %taddr, i64 2, <64 x i32> %stv64, i1 1)984 985  tail call void @llvm.nvvm.tcgen05.st.16x32bx2.x128(ptr addrspace(6) %taddr, i64 2, <128 x i32> %stv128, i1 1)986  ret void987}988