214 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 13 14declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)15declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)16declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)17declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)18 19define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {20; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg1(21; CHECK_PTX64: {22; CHECK_PTX64-NEXT: .reg .b32 %r<2>;23; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;24; CHECK_PTX64-EMPTY:25; CHECK_PTX64-NEXT: // %bb.0:26; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_cg1_param_0];27; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_cg1_param_1];28; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;29; CHECK_PTX64-NEXT: ret;30;31; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_cg1(32; CHECK_PTX64_SHARED32: {33; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;34; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;35; CHECK_PTX64_SHARED32-EMPTY:36; CHECK_PTX64_SHARED32-NEXT: // %bb.0:37; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_cg1_param_0];38; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_cg1_param_1];39; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;40; CHECK_PTX64_SHARED32-NEXT: ret;41 call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)42 ret void43}44 45define void @test_tcgen05_alloc_cg2(ptr %addr, i32 %ncols) {46; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg2(47; CHECK_PTX64: {48; CHECK_PTX64-NEXT: .reg .b32 %r<2>;49; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;50; CHECK_PTX64-EMPTY:51; CHECK_PTX64-NEXT: // %bb.0:52; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_cg2_param_0];53; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_cg2_param_1];54; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;55; CHECK_PTX64-NEXT: ret;56;57; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_cg2(58; CHECK_PTX64_SHARED32: {59; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;60; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;61; CHECK_PTX64_SHARED32-EMPTY:62; CHECK_PTX64_SHARED32-NEXT: // %bb.0:63; CHECK_PTX64_SHARED32-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_cg2_param_0];64; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_cg2_param_1];65; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;66; CHECK_PTX64_SHARED32-NEXT: ret;67 call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)68 ret void69}70 71define void @test_tcgen05_alloc_shared_cg1(ptr addrspace(3) %addr, i32 %ncols) {72; CHECK_PTX64-LABEL: test_tcgen05_alloc_shared_cg1(73; CHECK_PTX64: {74; CHECK_PTX64-NEXT: .reg .b32 %r<2>;75; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;76; CHECK_PTX64-EMPTY:77; CHECK_PTX64-NEXT: // %bb.0:78; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_shared_cg1_param_0];79; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg1_param_1];80; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%rd1], %r1;81; CHECK_PTX64-NEXT: ret;82;83; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_shared_cg1(84; CHECK_PTX64_SHARED32: {85; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;86; CHECK_PTX64_SHARED32-EMPTY:87; CHECK_PTX64_SHARED32-NEXT: // %bb.0:88; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg1_param_0];89; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_alloc_shared_cg1_param_1];90; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;91; CHECK_PTX64_SHARED32-NEXT: ret;92 call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)93 ret void94}95 96define void @test_tcgen05_alloc_shared_cg2(ptr addrspace(3) %addr, i32 %ncols) {97; CHECK_PTX64-LABEL: test_tcgen05_alloc_shared_cg2(98; CHECK_PTX64: {99; CHECK_PTX64-NEXT: .reg .b32 %r<2>;100; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;101; CHECK_PTX64-EMPTY:102; CHECK_PTX64-NEXT: // %bb.0:103; CHECK_PTX64-NEXT: ld.param.b64 %rd1, [test_tcgen05_alloc_shared_cg2_param_0];104; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg2_param_1];105; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%rd1], %r1;106; CHECK_PTX64-NEXT: ret;107;108; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_shared_cg2(109; CHECK_PTX64_SHARED32: {110; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;111; CHECK_PTX64_SHARED32-EMPTY:112; CHECK_PTX64_SHARED32-NEXT: // %bb.0:113; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg2_param_0];114; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_alloc_shared_cg2_param_1];115; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;116; CHECK_PTX64_SHARED32-NEXT: ret;117 call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)118 ret void119}120 121declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)122declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)123 124define void @test_tcgen05_dealloc_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {125; CHECK_PTX64-LABEL: test_tcgen05_dealloc_cg1(126; CHECK_PTX64: {127; CHECK_PTX64-NEXT: .reg .b32 %r<3>;128; CHECK_PTX64-EMPTY:129; CHECK_PTX64-NEXT: // %bb.0:130; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_dealloc_cg1_param_0];131; CHECK_PTX64-NEXT: ld.param.b32 %r2, [test_tcgen05_dealloc_cg1_param_1];132; CHECK_PTX64-NEXT: tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;133; CHECK_PTX64-NEXT: ret;134;135; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_cg1(136; CHECK_PTX64_SHARED32: {137; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;138; CHECK_PTX64_SHARED32-EMPTY:139; CHECK_PTX64_SHARED32-NEXT: // %bb.0:140; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_dealloc_cg1_param_0];141; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_dealloc_cg1_param_1];142; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;143; CHECK_PTX64_SHARED32-NEXT: ret;144 call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)145 ret void146}147 148define void @test_tcgen05_dealloc_cg2(ptr addrspace(6) %tmem_addr, i32 %ncols) {149; CHECK_PTX64-LABEL: test_tcgen05_dealloc_cg2(150; CHECK_PTX64: {151; CHECK_PTX64-NEXT: .reg .b32 %r<3>;152; CHECK_PTX64-EMPTY:153; CHECK_PTX64-NEXT: // %bb.0:154; CHECK_PTX64-NEXT: ld.param.b32 %r1, [test_tcgen05_dealloc_cg2_param_0];155; CHECK_PTX64-NEXT: ld.param.b32 %r2, [test_tcgen05_dealloc_cg2_param_1];156; CHECK_PTX64-NEXT: tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;157; CHECK_PTX64-NEXT: ret;158;159; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_cg2(160; CHECK_PTX64_SHARED32: {161; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;162; CHECK_PTX64_SHARED32-EMPTY:163; CHECK_PTX64_SHARED32-NEXT: // %bb.0:164; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r1, [test_tcgen05_dealloc_cg2_param_0];165; CHECK_PTX64_SHARED32-NEXT: ld.param.b32 %r2, [test_tcgen05_dealloc_cg2_param_1];166; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;167; CHECK_PTX64_SHARED32-NEXT: ret;168 call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)169 ret void170}171 172declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()173declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()174 175define void @test_tcgen05_relinquish_alloc_permit_cg1() {176; CHECK_PTX64-LABEL: test_tcgen05_relinquish_alloc_permit_cg1(177; CHECK_PTX64: {178; CHECK_PTX64-EMPTY:179; CHECK_PTX64-EMPTY:180; CHECK_PTX64-NEXT: // %bb.0:181; CHECK_PTX64-NEXT: tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;182; CHECK_PTX64-NEXT: ret;183;184; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_relinquish_alloc_permit_cg1(185; CHECK_PTX64_SHARED32: {186; CHECK_PTX64_SHARED32-EMPTY:187; CHECK_PTX64_SHARED32-EMPTY:188; CHECK_PTX64_SHARED32-NEXT: // %bb.0:189; CHECK_PTX64_SHARED32-NEXT: tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;190; CHECK_PTX64_SHARED32-NEXT: ret;191 call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()192 ret void193}194 195define void @test_tcgen05_relinquish_alloc_permit_cg2() {196; CHECK_PTX64-LABEL: test_tcgen05_relinquish_alloc_permit_cg2(197; CHECK_PTX64: {198; CHECK_PTX64-EMPTY:199; CHECK_PTX64-EMPTY:200; CHECK_PTX64-NEXT: // %bb.0:201; CHECK_PTX64-NEXT: tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;202; CHECK_PTX64-NEXT: ret;203;204; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_relinquish_alloc_permit_cg2(205; CHECK_PTX64_SHARED32: {206; CHECK_PTX64_SHARED32-EMPTY:207; CHECK_PTX64_SHARED32-EMPTY:208; CHECK_PTX64_SHARED32-NEXT: // %bb.0:209; CHECK_PTX64_SHARED32-NEXT: tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;210; CHECK_PTX64_SHARED32-NEXT: ret;211 call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()212 ret void213}214