218 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_PTX64 %s3; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s4; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s5; RUN: llc < %s -march=nvptx64 -mcpu=sm_100f -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s6; RUN: llc < %s -march=nvptx64 -mcpu=sm_110f -mattr=+ptx90 | FileCheck --check-prefixes=CHECK_PTX64 %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_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}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 13declare void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar_addr)14declare void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar_addr)15declare void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %bar_addr)16declare void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %bar_addr)17 18define void @test_tcgen05_commit_cg1(ptr %bar_addr) {19; CHECK_PTX64-LABEL: test_tcgen05_commit_cg1(20; CHECK_PTX64: {21; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;22; CHECK_PTX64-EMPTY:23; CHECK_PTX64-NEXT: // %bb.0:24; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg1_param_0];25; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%rd1];26; CHECK_PTX64-NEXT: ret;27;28; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_cg1(29; CHECK_PTX64_SHARED32: {30; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;31; CHECK_PTX64_SHARED32-EMPTY:32; CHECK_PTX64_SHARED32-NEXT: // %bb.0:33; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg1_param_0];34; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%rd1];35; CHECK_PTX64_SHARED32-NEXT: ret;36 call void @llvm.nvvm.tcgen05.commit.cg1(ptr %bar_addr)37 38 ret void39}40 41define void @test_tcgen05_commit_cg2(ptr %bar_addr) {42; CHECK_PTX64-LABEL: test_tcgen05_commit_cg2(43; CHECK_PTX64: {44; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;45; CHECK_PTX64-EMPTY:46; CHECK_PTX64-NEXT: // %bb.0:47; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg2_param_0];48; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%rd1];49; CHECK_PTX64-NEXT: ret;50;51; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_cg2(52; CHECK_PTX64_SHARED32: {53; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;54; CHECK_PTX64_SHARED32-EMPTY:55; CHECK_PTX64_SHARED32-NEXT: // %bb.0:56; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_cg2_param_0];57; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%rd1];58; CHECK_PTX64_SHARED32-NEXT: ret;59 call void @llvm.nvvm.tcgen05.commit.cg2(ptr %bar_addr)60 61 ret void62}63 64define void @test_tcgen05_commit_shared_cg1(ptr addrspace(3) %bar_addr) {65; CHECK_PTX64-LABEL: test_tcgen05_commit_shared_cg1(66; CHECK_PTX64: {67; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;68; CHECK_PTX64-EMPTY:69; CHECK_PTX64-NEXT: // %bb.0:70; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_shared_cg1_param_0];71; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%rd1];72; CHECK_PTX64-NEXT: ret;73;74; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_shared_cg1(75; CHECK_PTX64_SHARED32: {76; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;77; CHECK_PTX64_SHARED32-EMPTY:78; CHECK_PTX64_SHARED32-NEXT: // %bb.0:79; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_shared_cg1_param_0];80; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%r1];81; CHECK_PTX64_SHARED32-NEXT: ret;82 call void @llvm.nvvm.tcgen05.commit.shared.cg1(ptr addrspace(3) %bar_addr)83 84 ret void85}86 87define void @test_tcgen05_commit_shared_cg2(ptr addrspace(3) %bar_addr) {88; CHECK_PTX64-LABEL: test_tcgen05_commit_shared_cg2(89; CHECK_PTX64: {90; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;91; CHECK_PTX64-EMPTY:92; CHECK_PTX64-NEXT: // %bb.0:93; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_shared_cg2_param_0];94; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%rd1];95; CHECK_PTX64-NEXT: ret;96;97; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_shared_cg2(98; CHECK_PTX64_SHARED32: {99; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;100; CHECK_PTX64_SHARED32-EMPTY:101; CHECK_PTX64_SHARED32-NEXT: // %bb.0:102; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_shared_cg2_param_0];103; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64 [%r1];104; CHECK_PTX64_SHARED32-NEXT: ret;105 call void @llvm.nvvm.tcgen05.commit.shared.cg2(ptr addrspace(3) %bar_addr)106 107 ret void108}109 110declare void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar_addr, i16 %cta_mask)111declare void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar_addr, i16 %cta_mask)112declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask)113declare void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask)114 115define void @test_tcgen05_commit_mc_cg1(ptr %bar_addr, i16 %cta_mask) {116; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg1(117; CHECK_PTX64: {118; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;119; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;120; CHECK_PTX64-EMPTY:121; CHECK_PTX64-NEXT: // %bb.0:122; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg1_param_0];123; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg1_param_1];124; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;125; CHECK_PTX64-NEXT: ret;126;127; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_cg1(128; CHECK_PTX64_SHARED32: {129; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;130; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;131; CHECK_PTX64_SHARED32-EMPTY:132; CHECK_PTX64_SHARED32-NEXT: // %bb.0:133; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg1_param_0];134; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg1_param_1];135; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;136; CHECK_PTX64_SHARED32-NEXT: ret;137 call void @llvm.nvvm.tcgen05.commit.mc.cg1(ptr %bar_addr, i16 %cta_mask)138 ret void139}140 141define void @test_tcgen05_commit_mc_cg2(ptr %bar_addr, i16 %cta_mask) {142; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_cg2(143; CHECK_PTX64: {144; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;145; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;146; CHECK_PTX64-EMPTY:147; CHECK_PTX64-NEXT: // %bb.0:148; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg2_param_0];149; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg2_param_1];150; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;151; CHECK_PTX64-NEXT: ret;152;153; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_cg2(154; CHECK_PTX64_SHARED32: {155; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;156; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;157; CHECK_PTX64_SHARED32-EMPTY:158; CHECK_PTX64_SHARED32-NEXT: // %bb.0:159; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_cg2_param_0];160; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_cg2_param_1];161; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;162; CHECK_PTX64_SHARED32-NEXT: ret;163 call void @llvm.nvvm.tcgen05.commit.mc.cg2(ptr %bar_addr, i16 %cta_mask)164 ret void165}166 167define void @test_tcgen05_commit_mc_shared_cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask) {168; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_shared_cg1(169; CHECK_PTX64: {170; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;171; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;172; CHECK_PTX64-EMPTY:173; CHECK_PTX64-NEXT: // %bb.0:174; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_shared_cg1_param_0];175; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg1_param_1];176; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;177; CHECK_PTX64-NEXT: ret;178;179; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_shared_cg1(180; CHECK_PTX64_SHARED32: {181; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;182; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;183; CHECK_PTX64_SHARED32-EMPTY:184; CHECK_PTX64_SHARED32-NEXT: // %bb.0:185; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg1_param_0];186; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg1_param_1];187; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%r1], %rs1;188; CHECK_PTX64_SHARED32-NEXT: ret;189 call void @llvm.nvvm.tcgen05.commit.mc.shared.cg1(ptr addrspace(3) %bar_addr, i16 %cta_mask)190 ret void191}192 193define void @test_tcgen05_commit_mc_shared_cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask) {194; CHECK_PTX64-LABEL: test_tcgen05_commit_mc_shared_cg2(195; CHECK_PTX64: {196; CHECK_PTX64-NEXT: .reg .b16 %rs<2>;197; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;198; CHECK_PTX64-EMPTY:199; CHECK_PTX64-NEXT: // %bb.0:200; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_commit_mc_shared_cg2_param_0];201; CHECK_PTX64-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg2_param_1];202; CHECK_PTX64-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%rd1], %rs1;203; CHECK_PTX64-NEXT: ret;204;205; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_commit_mc_shared_cg2(206; CHECK_PTX64_SHARED32: {207; CHECK_PTX64_SHARED32-NEXT: .reg .b16 %rs<2>;208; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;209; CHECK_PTX64_SHARED32-EMPTY:210; CHECK_PTX64_SHARED32-NEXT: // %bb.0:211; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_commit_mc_shared_cg2_param_0];212; CHECK_PTX64_SHARED32-NEXT: ld.param.b16 %rs1, [test_tcgen05_commit_mc_shared_cg2_param_1];213; CHECK_PTX64_SHARED32-NEXT: tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [%r1], %rs1;214; CHECK_PTX64_SHARED32-NEXT: ret;215 call void @llvm.nvvm.tcgen05.commit.mc.shared.cg2(ptr addrspace(3) %bar_addr, i16 %cta_mask)216 ret void217}218