589 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK %s3; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s4; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK %s5; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK %s6; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}7; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}8; RUN: %if ptxas-sm_100f && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | %ptxas-verify -arch=sm_100f %}9; RUN: %if ptxas-sm_110f && ptxas-isa-9.0 %{ llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | %ptxas-verify -arch=sm_110f %}10 11define void @test_tcgen05_cp_64x128_v1_cg1(ptr addrspace(6) %addr, i64 %sdesc) {12; CHECK-LABEL: test_tcgen05_cp_64x128_v1_cg1(13; CHECK: {14; CHECK-NEXT: .reg .b32 %r<2>;15; CHECK-NEXT: .reg .b64 %rd<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_cg1_param_0];19; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_cg1_param_1];20; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::02_13 [%r1], %rd1;21; CHECK-NEXT: ret;22 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.cg1(ptr addrspace(6) %addr, i64 %sdesc)23 24 ret void25}26 27define void @test_tcgen05_cp_64x128_v1_cg2(ptr addrspace(6) %addr, i64 %sdesc) {28; CHECK-LABEL: test_tcgen05_cp_64x128_v1_cg2(29; CHECK: {30; CHECK-NEXT: .reg .b32 %r<2>;31; CHECK-NEXT: .reg .b64 %rd<2>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_cg2_param_0];35; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_cg2_param_1];36; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::02_13 [%r1], %rd1;37; CHECK-NEXT: ret;38 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.cg2(ptr addrspace(6) %addr, i64 %sdesc)39 40 ret void41}42 43define void @test_tcgen05_cp_64x128_v2_cg1(ptr addrspace(6) %addr, i64 %sdesc) {44; CHECK-LABEL: test_tcgen05_cp_64x128_v2_cg1(45; CHECK: {46; CHECK-NEXT: .reg .b32 %r<2>;47; CHECK-NEXT: .reg .b64 %rd<2>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_cg1_param_0];51; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_cg1_param_1];52; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::01_23 [%r1], %rd1;53; CHECK-NEXT: ret;54 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.cg1(ptr addrspace(6) %addr, i64 %sdesc)55 56 ret void57}58 59define void @test_tcgen05_cp_64x128_v2_cg2(ptr addrspace(6) %addr, i64 %sdesc) {60; CHECK-LABEL: test_tcgen05_cp_64x128_v2_cg2(61; CHECK: {62; CHECK-NEXT: .reg .b32 %r<2>;63; CHECK-NEXT: .reg .b64 %rd<2>;64; CHECK-EMPTY:65; CHECK-NEXT: // %bb.0:66; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_cg2_param_0];67; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_cg2_param_1];68; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::01_23 [%r1], %rd1;69; CHECK-NEXT: ret;70 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.cg2(ptr addrspace(6) %addr, i64 %sdesc)71 72 ret void73}74 75define void @test_tcgen05_cp_32x128_cg1(ptr addrspace(6) %addr, i64 %sdesc) {76; CHECK-LABEL: test_tcgen05_cp_32x128_cg1(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<2>;79; CHECK-NEXT: .reg .b64 %rd<2>;80; CHECK-EMPTY:81; CHECK-NEXT: // %bb.0:82; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_cg1_param_0];83; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_cg1_param_1];84; CHECK-NEXT: tcgen05.cp.cta_group::1.32x128b.warpx4 [%r1], %rd1;85; CHECK-NEXT: ret;86 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.cg1(ptr addrspace(6) %addr, i64 %sdesc)87 88 ret void89}90 91define void @test_tcgen05_cp_32x128_cg2(ptr addrspace(6) %addr, i64 %sdesc) {92; CHECK-LABEL: test_tcgen05_cp_32x128_cg2(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<2>;95; CHECK-NEXT: .reg .b64 %rd<2>;96; CHECK-EMPTY:97; CHECK-NEXT: // %bb.0:98; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_cg2_param_0];99; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_cg2_param_1];100; CHECK-NEXT: tcgen05.cp.cta_group::2.32x128b.warpx4 [%r1], %rd1;101; CHECK-NEXT: ret;102 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.cg2(ptr addrspace(6) %addr, i64 %sdesc)103 104 ret void105}106 107 108define void @test_tcgen05_cp_128x128b_cg1(ptr addrspace(6) %addr, i64 %sdesc) {109; CHECK-LABEL: test_tcgen05_cp_128x128b_cg1(110; CHECK: {111; CHECK-NEXT: .reg .b32 %r<2>;112; CHECK-NEXT: .reg .b64 %rd<2>;113; CHECK-EMPTY:114; CHECK-NEXT: // %bb.0:115; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_cg1_param_0];116; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_cg1_param_1];117; CHECK-NEXT: tcgen05.cp.cta_group::1.128x128b [%r1], %rd1;118; CHECK-NEXT: ret;119 call void @llvm.nvvm.tcgen05.cp.128x128b.cg1(ptr addrspace(6) %addr, i64 %sdesc)120 121 ret void122}123 124define void @test_tcgen05_cp_128x128b_cg2(ptr addrspace(6) %addr, i64 %sdesc) {125; CHECK-LABEL: test_tcgen05_cp_128x128b_cg2(126; CHECK: {127; CHECK-NEXT: .reg .b32 %r<2>;128; CHECK-NEXT: .reg .b64 %rd<2>;129; CHECK-EMPTY:130; CHECK-NEXT: // %bb.0:131; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_cg2_param_0];132; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_cg2_param_1];133; CHECK-NEXT: tcgen05.cp.cta_group::2.128x128b [%r1], %rd1;134; CHECK-NEXT: ret;135 call void @llvm.nvvm.tcgen05.cp.128x128b.cg2(ptr addrspace(6) %addr, i64 %sdesc)136 137 ret void138}139 140define void @test_tcgen05_cp_128x256b_cg1(ptr addrspace(6) %addr, i64 %sdesc) {141; CHECK-LABEL: test_tcgen05_cp_128x256b_cg1(142; CHECK: {143; CHECK-NEXT: .reg .b32 %r<2>;144; CHECK-NEXT: .reg .b64 %rd<2>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_cg1_param_0];148; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_cg1_param_1];149; CHECK-NEXT: tcgen05.cp.cta_group::1.128x256b [%r1], %rd1;150; CHECK-NEXT: ret;151 call void @llvm.nvvm.tcgen05.cp.128x256b.cg1(ptr addrspace(6) %addr, i64 %sdesc)152 153 ret void154}155 156define void @test_tcgen05_cp_128x256b_cg2(ptr addrspace(6) %addr, i64 %sdesc) {157; CHECK-LABEL: test_tcgen05_cp_128x256b_cg2(158; CHECK: {159; CHECK-NEXT: .reg .b32 %r<2>;160; CHECK-NEXT: .reg .b64 %rd<2>;161; CHECK-EMPTY:162; CHECK-NEXT: // %bb.0:163; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_cg2_param_0];164; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_cg2_param_1];165; CHECK-NEXT: tcgen05.cp.cta_group::2.128x256b [%r1], %rd1;166; CHECK-NEXT: ret;167 call void @llvm.nvvm.tcgen05.cp.128x256b.cg2(ptr addrspace(6) %addr, i64 %sdesc)168 169 ret void170}171 172define void @test_tcgen05_cp_4x256b_cg1(ptr addrspace(6) %addr, i64 %sdesc) {173; CHECK-LABEL: test_tcgen05_cp_4x256b_cg1(174; CHECK: {175; CHECK-NEXT: .reg .b32 %r<2>;176; CHECK-NEXT: .reg .b64 %rd<2>;177; CHECK-EMPTY:178; CHECK-NEXT: // %bb.0:179; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_cg1_param_0];180; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_cg1_param_1];181; CHECK-NEXT: tcgen05.cp.cta_group::1.4x256b [%r1], %rd1;182; CHECK-NEXT: ret;183 call void @llvm.nvvm.tcgen05.cp.4x256b.cg1(ptr addrspace(6) %addr, i64 %sdesc)184 185 ret void186}187 188define void @test_tcgen05_cp_4x256b_cg2(ptr addrspace(6) %addr, i64 %sdesc) {189; CHECK-LABEL: test_tcgen05_cp_4x256b_cg2(190; CHECK: {191; CHECK-NEXT: .reg .b32 %r<2>;192; CHECK-NEXT: .reg .b64 %rd<2>;193; CHECK-EMPTY:194; CHECK-NEXT: // %bb.0:195; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_cg2_param_0];196; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_cg2_param_1];197; CHECK-NEXT: tcgen05.cp.cta_group::2.4x256b [%r1], %rd1;198; CHECK-NEXT: ret;199 call void @llvm.nvvm.tcgen05.cp.4x256b.cg2(ptr addrspace(6) %addr, i64 %sdesc)200 201 ret void202}203 204; With src_fmt as b6x16_p32205define void @test_tcgen05_cp_128x256b_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {206; CHECK-LABEL: test_tcgen05_cp_128x256b_b6x16_p32_cg1(207; CHECK: {208; CHECK-NEXT: .reg .b32 %r<2>;209; CHECK-NEXT: .reg .b64 %rd<2>;210; CHECK-EMPTY:211; CHECK-NEXT: // %bb.0:212; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_b6x16_p32_cg1_param_0];213; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_b6x16_p32_cg1_param_1];214; CHECK-NEXT: tcgen05.cp.cta_group::1.128x256b.b8x16.b6x16_p32 [%r1], %rd1;215; CHECK-NEXT: ret;216 call void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)217 218 ret void219}220 221define void @test_tcgen05_cp_128x256b_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {222; CHECK-LABEL: test_tcgen05_cp_128x256b_b6x16_p32_cg2(223; CHECK: {224; CHECK-NEXT: .reg .b32 %r<2>;225; CHECK-NEXT: .reg .b64 %rd<2>;226; CHECK-EMPTY:227; CHECK-NEXT: // %bb.0:228; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_b6x16_p32_cg2_param_0];229; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_b6x16_p32_cg2_param_1];230; CHECK-NEXT: tcgen05.cp.cta_group::2.128x256b.b8x16.b6x16_p32 [%r1], %rd1;231; CHECK-NEXT: ret;232 call void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)233 234 ret void235}236 237define void @test_tcgen05_cp_4x256b_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {238; CHECK-LABEL: test_tcgen05_cp_4x256b_b6x16_p32_cg1(239; CHECK: {240; CHECK-NEXT: .reg .b32 %r<2>;241; CHECK-NEXT: .reg .b64 %rd<2>;242; CHECK-EMPTY:243; CHECK-NEXT: // %bb.0:244; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_b6x16_p32_cg1_param_0];245; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_b6x16_p32_cg1_param_1];246; CHECK-NEXT: tcgen05.cp.cta_group::1.4x256b.b8x16.b6x16_p32 [%r1], %rd1;247; CHECK-NEXT: ret;248 call void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)249 250 ret void251}252 253define void @test_tcgen05_cp_4x256b_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {254; CHECK-LABEL: test_tcgen05_cp_4x256b_b6x16_p32_cg2(255; CHECK: {256; CHECK-NEXT: .reg .b32 %r<2>;257; CHECK-NEXT: .reg .b64 %rd<2>;258; CHECK-EMPTY:259; CHECK-NEXT: // %bb.0:260; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_b6x16_p32_cg2_param_0];261; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_b6x16_p32_cg2_param_1];262; CHECK-NEXT: tcgen05.cp.cta_group::2.4x256b.b8x16.b6x16_p32 [%r1], %rd1;263; CHECK-NEXT: ret;264 call void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)265 266 ret void267}268 269define void @test_tcgen05_cp_128x128b_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {270; CHECK-LABEL: test_tcgen05_cp_128x128b_b6x16_p32_cg1(271; CHECK: {272; CHECK-NEXT: .reg .b32 %r<2>;273; CHECK-NEXT: .reg .b64 %rd<2>;274; CHECK-EMPTY:275; CHECK-NEXT: // %bb.0:276; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_b6x16_p32_cg1_param_0];277; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_b6x16_p32_cg1_param_1];278; CHECK-NEXT: tcgen05.cp.cta_group::1.128x128b.b8x16.b6x16_p32 [%r1], %rd1;279; CHECK-NEXT: ret;280 call void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)281 282 ret void283}284 285define void @test_tcgen05_cp_128x128b_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {286; CHECK-LABEL: test_tcgen05_cp_128x128b_b6x16_p32_cg2(287; CHECK: {288; CHECK-NEXT: .reg .b32 %r<2>;289; CHECK-NEXT: .reg .b64 %rd<2>;290; CHECK-EMPTY:291; CHECK-NEXT: // %bb.0:292; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_b6x16_p32_cg2_param_0];293; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_b6x16_p32_cg2_param_1];294; CHECK-NEXT: tcgen05.cp.cta_group::2.128x128b.b8x16.b6x16_p32 [%r1], %rd1;295; CHECK-NEXT: ret;296 call void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)297 298 ret void299}300 301define void @test_tcgen05_cp_64x128_v1_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {302; CHECK-LABEL: test_tcgen05_cp_64x128_v1_b6x16_p32_cg1(303; CHECK: {304; CHECK-NEXT: .reg .b32 %r<2>;305; CHECK-NEXT: .reg .b64 %rd<2>;306; CHECK-EMPTY:307; CHECK-NEXT: // %bb.0:308; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_b6x16_p32_cg1_param_0];309; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_b6x16_p32_cg1_param_1];310; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::02_13.b8x16.b6x16_p32 [%r1], %rd1;311; CHECK-NEXT: ret;312 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)313 314 ret void315}316 317define void @test_tcgen05_cp_64x128_v1_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {318; CHECK-LABEL: test_tcgen05_cp_64x128_v1_b6x16_p32_cg2(319; CHECK: {320; CHECK-NEXT: .reg .b32 %r<2>;321; CHECK-NEXT: .reg .b64 %rd<2>;322; CHECK-EMPTY:323; CHECK-NEXT: // %bb.0:324; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_b6x16_p32_cg2_param_0];325; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_b6x16_p32_cg2_param_1];326; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::02_13.b8x16.b6x16_p32 [%r1], %rd1;327; CHECK-NEXT: ret;328 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)329 330 ret void331}332 333define void @test_tcgen05_cp_64x128_v2_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {334; CHECK-LABEL: test_tcgen05_cp_64x128_v2_b6x16_p32_cg1(335; CHECK: {336; CHECK-NEXT: .reg .b32 %r<2>;337; CHECK-NEXT: .reg .b64 %rd<2>;338; CHECK-EMPTY:339; CHECK-NEXT: // %bb.0:340; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_b6x16_p32_cg1_param_0];341; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_b6x16_p32_cg1_param_1];342; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::01_23.b8x16.b6x16_p32 [%r1], %rd1;343; CHECK-NEXT: ret;344 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)345 346 ret void347}348 349define void @test_tcgen05_cp_64x128_v2_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {350; CHECK-LABEL: test_tcgen05_cp_64x128_v2_b6x16_p32_cg2(351; CHECK: {352; CHECK-NEXT: .reg .b32 %r<2>;353; CHECK-NEXT: .reg .b64 %rd<2>;354; CHECK-EMPTY:355; CHECK-NEXT: // %bb.0:356; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_b6x16_p32_cg2_param_0];357; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_b6x16_p32_cg2_param_1];358; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::01_23.b8x16.b6x16_p32 [%r1], %rd1;359; CHECK-NEXT: ret;360 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)361 362 ret void363}364 365define void @test_tcgen05_cp_32x128_b6x16_p32_cg1(ptr addrspace(6) %addr, i64 %sdesc) {366; CHECK-LABEL: test_tcgen05_cp_32x128_b6x16_p32_cg1(367; CHECK: {368; CHECK-NEXT: .reg .b32 %r<2>;369; CHECK-NEXT: .reg .b64 %rd<2>;370; CHECK-EMPTY:371; CHECK-NEXT: // %bb.0:372; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_b6x16_p32_cg1_param_0];373; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_b6x16_p32_cg1_param_1];374; CHECK-NEXT: tcgen05.cp.cta_group::1.32x128b.warpx4.b8x16.b6x16_p32 [%r1], %rd1;375; CHECK-NEXT: ret;376 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.cg1(ptr addrspace(6) %addr, i64 %sdesc)377 378 ret void379}380 381define void @test_tcgen05_cp_32x128_b6x16_p32_cg2(ptr addrspace(6) %addr, i64 %sdesc) {382; CHECK-LABEL: test_tcgen05_cp_32x128_b6x16_p32_cg2(383; CHECK: {384; CHECK-NEXT: .reg .b32 %r<2>;385; CHECK-NEXT: .reg .b64 %rd<2>;386; CHECK-EMPTY:387; CHECK-NEXT: // %bb.0:388; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_b6x16_p32_cg2_param_0];389; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_b6x16_p32_cg2_param_1];390; CHECK-NEXT: tcgen05.cp.cta_group::2.32x128b.warpx4.b8x16.b6x16_p32 [%r1], %rd1;391; CHECK-NEXT: ret;392 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.cg2(ptr addrspace(6) %addr, i64 %sdesc)393 394 ret void395}396 397; With src_fmt as b4x16_p64398define void @test_tcgen05_cp_128x256b_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {399; CHECK-LABEL: test_tcgen05_cp_128x256b_b4x16_p64_cg1(400; CHECK: {401; CHECK-NEXT: .reg .b32 %r<2>;402; CHECK-NEXT: .reg .b64 %rd<2>;403; CHECK-EMPTY:404; CHECK-NEXT: // %bb.0:405; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_b4x16_p64_cg1_param_0];406; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_b4x16_p64_cg1_param_1];407; CHECK-NEXT: tcgen05.cp.cta_group::1.128x256b.b8x16.b4x16_p64 [%r1], %rd1;408; CHECK-NEXT: ret;409 call void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)410 411 ret void412}413 414define void @test_tcgen05_cp_128x256b_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {415; CHECK-LABEL: test_tcgen05_cp_128x256b_b4x16_p64_cg2(416; CHECK: {417; CHECK-NEXT: .reg .b32 %r<2>;418; CHECK-NEXT: .reg .b64 %rd<2>;419; CHECK-EMPTY:420; CHECK-NEXT: // %bb.0:421; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x256b_b4x16_p64_cg2_param_0];422; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x256b_b4x16_p64_cg2_param_1];423; CHECK-NEXT: tcgen05.cp.cta_group::2.128x256b.b8x16.b4x16_p64 [%r1], %rd1;424; CHECK-NEXT: ret;425 call void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)426 427 ret void428}429 430define void @test_tcgen05_cp_4x256b_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {431; CHECK-LABEL: test_tcgen05_cp_4x256b_b4x16_p64_cg1(432; CHECK: {433; CHECK-NEXT: .reg .b32 %r<2>;434; CHECK-NEXT: .reg .b64 %rd<2>;435; CHECK-EMPTY:436; CHECK-NEXT: // %bb.0:437; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_b4x16_p64_cg1_param_0];438; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_b4x16_p64_cg1_param_1];439; CHECK-NEXT: tcgen05.cp.cta_group::1.4x256b.b8x16.b4x16_p64 [%r1], %rd1;440; CHECK-NEXT: ret;441 call void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)442 443 ret void444}445 446define void @test_tcgen05_cp_4x256b_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {447; CHECK-LABEL: test_tcgen05_cp_4x256b_b4x16_p64_cg2(448; CHECK: {449; CHECK-NEXT: .reg .b32 %r<2>;450; CHECK-NEXT: .reg .b64 %rd<2>;451; CHECK-EMPTY:452; CHECK-NEXT: // %bb.0:453; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_4x256b_b4x16_p64_cg2_param_0];454; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_4x256b_b4x16_p64_cg2_param_1];455; CHECK-NEXT: tcgen05.cp.cta_group::2.4x256b.b8x16.b4x16_p64 [%r1], %rd1;456; CHECK-NEXT: ret;457 call void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)458 459 ret void460}461 462define void @test_tcgen05_cp_128x128b_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {463; CHECK-LABEL: test_tcgen05_cp_128x128b_b4x16_p64_cg1(464; CHECK: {465; CHECK-NEXT: .reg .b32 %r<2>;466; CHECK-NEXT: .reg .b64 %rd<2>;467; CHECK-EMPTY:468; CHECK-NEXT: // %bb.0:469; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_b4x16_p64_cg1_param_0];470; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_b4x16_p64_cg1_param_1];471; CHECK-NEXT: tcgen05.cp.cta_group::1.128x128b.b8x16.b4x16_p64 [%r1], %rd1;472; CHECK-NEXT: ret;473 call void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)474 475 ret void476}477 478define void @test_tcgen05_cp_128x128b_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {479; CHECK-LABEL: test_tcgen05_cp_128x128b_b4x16_p64_cg2(480; CHECK: {481; CHECK-NEXT: .reg .b32 %r<2>;482; CHECK-NEXT: .reg .b64 %rd<2>;483; CHECK-EMPTY:484; CHECK-NEXT: // %bb.0:485; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_128x128b_b4x16_p64_cg2_param_0];486; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_128x128b_b4x16_p64_cg2_param_1];487; CHECK-NEXT: tcgen05.cp.cta_group::2.128x128b.b8x16.b4x16_p64 [%r1], %rd1;488; CHECK-NEXT: ret;489 call void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)490 491 ret void492}493 494define void @test_tcgen05_cp_64x128_v1_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {495; CHECK-LABEL: test_tcgen05_cp_64x128_v1_b4x16_p64_cg1(496; CHECK: {497; CHECK-NEXT: .reg .b32 %r<2>;498; CHECK-NEXT: .reg .b64 %rd<2>;499; CHECK-EMPTY:500; CHECK-NEXT: // %bb.0:501; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_b4x16_p64_cg1_param_0];502; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_b4x16_p64_cg1_param_1];503; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::02_13.b8x16.b4x16_p64 [%r1], %rd1;504; CHECK-NEXT: ret;505 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)506 507 ret void508}509 510define void @test_tcgen05_cp_64x128_v1_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {511; CHECK-LABEL: test_tcgen05_cp_64x128_v1_b4x16_p64_cg2(512; CHECK: {513; CHECK-NEXT: .reg .b32 %r<2>;514; CHECK-NEXT: .reg .b64 %rd<2>;515; CHECK-EMPTY:516; CHECK-NEXT: // %bb.0:517; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v1_b4x16_p64_cg2_param_0];518; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v1_b4x16_p64_cg2_param_1];519; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::02_13.b8x16.b4x16_p64 [%r1], %rd1;520; CHECK-NEXT: ret;521 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)522 523 ret void524}525 526define void @test_tcgen05_cp_64x128_v2_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {527; CHECK-LABEL: test_tcgen05_cp_64x128_v2_b4x16_p64_cg1(528; CHECK: {529; CHECK-NEXT: .reg .b32 %r<2>;530; CHECK-NEXT: .reg .b64 %rd<2>;531; CHECK-EMPTY:532; CHECK-NEXT: // %bb.0:533; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_b4x16_p64_cg1_param_0];534; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_b4x16_p64_cg1_param_1];535; CHECK-NEXT: tcgen05.cp.cta_group::1.64x128b.warpx2::01_23.b8x16.b4x16_p64 [%r1], %rd1;536; CHECK-NEXT: ret;537 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)538 539 ret void540}541 542define void @test_tcgen05_cp_64x128_v2_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {543; CHECK-LABEL: test_tcgen05_cp_64x128_v2_b4x16_p64_cg2(544; CHECK: {545; CHECK-NEXT: .reg .b32 %r<2>;546; CHECK-NEXT: .reg .b64 %rd<2>;547; CHECK-EMPTY:548; CHECK-NEXT: // %bb.0:549; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_64x128_v2_b4x16_p64_cg2_param_0];550; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_64x128_v2_b4x16_p64_cg2_param_1];551; CHECK-NEXT: tcgen05.cp.cta_group::2.64x128b.warpx2::01_23.b8x16.b4x16_p64 [%r1], %rd1;552; CHECK-NEXT: ret;553 call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)554 555 ret void556}557 558define void @test_tcgen05_cp_32x128_b4x16_p64_cg1(ptr addrspace(6) %addr, i64 %sdesc) {559; CHECK-LABEL: test_tcgen05_cp_32x128_b4x16_p64_cg1(560; CHECK: {561; CHECK-NEXT: .reg .b32 %r<2>;562; CHECK-NEXT: .reg .b64 %rd<2>;563; CHECK-EMPTY:564; CHECK-NEXT: // %bb.0:565; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_b4x16_p64_cg1_param_0];566; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_b4x16_p64_cg1_param_1];567; CHECK-NEXT: tcgen05.cp.cta_group::1.32x128b.warpx4.b8x16.b4x16_p64 [%r1], %rd1;568; CHECK-NEXT: ret;569 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.cg1(ptr addrspace(6) %addr, i64 %sdesc)570 571 ret void572}573 574define void @test_tcgen05_cp_32x128_b4x16_p64_cg2(ptr addrspace(6) %addr, i64 %sdesc) {575; CHECK-LABEL: test_tcgen05_cp_32x128_b4x16_p64_cg2(576; CHECK: {577; CHECK-NEXT: .reg .b32 %r<2>;578; CHECK-NEXT: .reg .b64 %rd<2>;579; CHECK-EMPTY:580; CHECK-NEXT: // %bb.0:581; CHECK-NEXT: ld.param.b32 %r1, [test_tcgen05_cp_32x128_b4x16_p64_cg2_param_0];582; CHECK-NEXT: ld.param.b64 %rd1, [test_tcgen05_cp_32x128_b4x16_p64_cg2_param_1];583; CHECK-NEXT: tcgen05.cp.cta_group::2.32x128b.warpx4.b8x16.b4x16_p64 [%r1], %rd1;584; CHECK-NEXT: ret;585 call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.cg2(ptr addrspace(6) %addr, i64 %sdesc)586 587 ret void588}589