88 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 62; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}5; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}6 7declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cta(ptr addrspace(3), i32)8declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cta(ptr addrspace(3), i32)9declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cta(ptr addrspace(3), i32)10declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cta(ptr addrspace(3), i32)11 12declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cluster(ptr addrspace(7), i32)13declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cluster(ptr addrspace(7), i32)14declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cluster(ptr addrspace(7), i32)15declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cluster(ptr addrspace(7), i32)16 17define void @test_mbarrier_tx_space_cta(ptr addrspace(3) %mbar, i32 %tx_count) {18; CHECK-PTX64-LABEL: test_mbarrier_tx_space_cta(19; CHECK-PTX64: {20; CHECK-PTX64-NEXT: .reg .b32 %r<2>;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_mbarrier_tx_space_cta_param_0];25; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_tx_space_cta_param_1];26; CHECK-PTX64-NEXT: mbarrier.expect_tx.relaxed.cta.shared.b64 [%rd1], %r1;27; CHECK-PTX64-NEXT: mbarrier.expect_tx.relaxed.cluster.shared.b64 [%rd1], %r1;28; CHECK-PTX64-NEXT: mbarrier.complete_tx.relaxed.cta.shared.b64 [%rd1], %r1;29; CHECK-PTX64-NEXT: mbarrier.complete_tx.relaxed.cluster.shared.b64 [%rd1], %r1;30; CHECK-PTX64-NEXT: ret;31;32; CHECK-PTX-SHARED32-LABEL: test_mbarrier_tx_space_cta(33; CHECK-PTX-SHARED32: {34; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;35; CHECK-PTX-SHARED32-EMPTY:36; CHECK-PTX-SHARED32-NEXT: // %bb.0:37; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_tx_space_cta_param_0];38; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_tx_space_cta_param_1];39; CHECK-PTX-SHARED32-NEXT: mbarrier.expect_tx.relaxed.cta.shared.b64 [%r1], %r2;40; CHECK-PTX-SHARED32-NEXT: mbarrier.expect_tx.relaxed.cluster.shared.b64 [%r1], %r2;41; CHECK-PTX-SHARED32-NEXT: mbarrier.complete_tx.relaxed.cta.shared.b64 [%r1], %r2;42; CHECK-PTX-SHARED32-NEXT: mbarrier.complete_tx.relaxed.cluster.shared.b64 [%r1], %r2;43; CHECK-PTX-SHARED32-NEXT: ret;44 tail call void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx_count)45 tail call void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx_count)46 47 tail call void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx_count)48 tail call void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx_count)49 50 ret void51}52 53define void @test_mbarrier_tx_space_cluster(ptr addrspace(7) %mbar, i32 %tx_count) {54; CHECK-PTX64-LABEL: test_mbarrier_tx_space_cluster(55; CHECK-PTX64: {56; CHECK-PTX64-NEXT: .reg .b32 %r<2>;57; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;58; CHECK-PTX64-EMPTY:59; CHECK-PTX64-NEXT: // %bb.0:60; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_mbarrier_tx_space_cluster_param_0];61; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_tx_space_cluster_param_1];62; CHECK-PTX64-NEXT: mbarrier.expect_tx.relaxed.cta.shared::cluster.b64 [%rd1], %r1;63; CHECK-PTX64-NEXT: mbarrier.expect_tx.relaxed.cluster.shared::cluster.b64 [%rd1], %r1;64; CHECK-PTX64-NEXT: mbarrier.complete_tx.relaxed.cta.shared::cluster.b64 [%rd1], %r1;65; CHECK-PTX64-NEXT: mbarrier.complete_tx.relaxed.cluster.shared::cluster.b64 [%rd1], %r1;66; CHECK-PTX64-NEXT: ret;67;68; CHECK-PTX-SHARED32-LABEL: test_mbarrier_tx_space_cluster(69; CHECK-PTX-SHARED32: {70; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;71; CHECK-PTX-SHARED32-EMPTY:72; CHECK-PTX-SHARED32-NEXT: // %bb.0:73; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_tx_space_cluster_param_0];74; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_tx_space_cluster_param_1];75; CHECK-PTX-SHARED32-NEXT: mbarrier.expect_tx.relaxed.cta.shared::cluster.b64 [%r1], %r2;76; CHECK-PTX-SHARED32-NEXT: mbarrier.expect_tx.relaxed.cluster.shared::cluster.b64 [%r1], %r2;77; CHECK-PTX-SHARED32-NEXT: mbarrier.complete_tx.relaxed.cta.shared::cluster.b64 [%r1], %r2;78; CHECK-PTX-SHARED32-NEXT: mbarrier.complete_tx.relaxed.cluster.shared::cluster.b64 [%r1], %r2;79; CHECK-PTX-SHARED32-NEXT: ret;80 tail call void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx_count)81 tail call void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx_count)82 83 tail call void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx_count)84 tail call void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx_count)85 86 ret void87}88