166 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 7; === space_cta (addrspace 3) ===8define void @test_mbarrier_arrive_scope_cta_space_cta(ptr addrspace(3) %mbar, i32 %tx) {9; CHECK-PTX64-LABEL: test_mbarrier_arrive_scope_cta_space_cta(10; CHECK-PTX64: {11; CHECK-PTX64-NEXT: .reg .b32 %r<2>;12; CHECK-PTX64-NEXT: .reg .b64 %rd<6>;13; CHECK-PTX64-EMPTY:14; CHECK-PTX64-NEXT: // %bb.0:15; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_mbarrier_arrive_scope_cta_space_cta_param_0];16; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cta_space_cta_param_1];17; CHECK-PTX64-NEXT: mbarrier.arrive.shared.b64 %rd2, [%rd1], %r1;18; CHECK-PTX64-NEXT: mbarrier.arrive.expect_tx.shared.b64 %rd3, [%rd1], %r1;19; CHECK-PTX64-NEXT: mbarrier.arrive_drop.shared.b64 %rd4, [%rd1], %r1;20; CHECK-PTX64-NEXT: mbarrier.arrive_drop.expect_tx.shared.b64 %rd5, [%rd1], %r1;21; CHECK-PTX64-NEXT: ret;22;23; CHECK-PTX-SHARED32-LABEL: test_mbarrier_arrive_scope_cta_space_cta(24; CHECK-PTX-SHARED32: {25; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;26; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<5>;27; CHECK-PTX-SHARED32-EMPTY:28; CHECK-PTX-SHARED32-NEXT: // %bb.0:29; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cta_space_cta_param_0];30; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_arrive_scope_cta_space_cta_param_1];31; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.shared.b64 %rd1, [%r1], %r2;32; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.expect_tx.shared.b64 %rd2, [%r1], %r2;33; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.shared.b64 %rd3, [%r1], %r2;34; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.expect_tx.shared.b64 %rd4, [%r1], %r2;35; CHECK-PTX-SHARED32-NEXT: ret;36 %r1 = call i64 @llvm.nvvm.mbarrier.arrive.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx)37 %r2 = call i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx)38 %r3 = call i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx)39 %r4 = call i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %tx)40 ret void41}42 43define void @test_mbarrier_arrive_scope_cluster_space_cta(ptr addrspace(3) %mbar, i32 %tx) {44; CHECK-PTX64-LABEL: test_mbarrier_arrive_scope_cluster_space_cta(45; CHECK-PTX64: {46; CHECK-PTX64-NEXT: .reg .b32 %r<2>;47; CHECK-PTX64-NEXT: .reg .b64 %rd<6>;48; CHECK-PTX64-EMPTY:49; CHECK-PTX64-NEXT: // %bb.0:50; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_mbarrier_arrive_scope_cluster_space_cta_param_0];51; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cluster_space_cta_param_1];52; CHECK-PTX64-NEXT: mbarrier.arrive.release.cluster.shared.b64 %rd2, [%rd1], %r1;53; CHECK-PTX64-NEXT: mbarrier.arrive.expect_tx.release.cluster.shared.b64 %rd3, [%rd1], %r1;54; CHECK-PTX64-NEXT: mbarrier.arrive_drop.release.cluster.shared.b64 %rd4, [%rd1], %r1;55; CHECK-PTX64-NEXT: mbarrier.arrive_drop.expect_tx.release.cluster.shared.b64 %rd5, [%rd1], %r1;56; CHECK-PTX64-NEXT: ret;57;58; CHECK-PTX-SHARED32-LABEL: test_mbarrier_arrive_scope_cluster_space_cta(59; CHECK-PTX-SHARED32: {60; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;61; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<5>;62; CHECK-PTX-SHARED32-EMPTY:63; CHECK-PTX-SHARED32-NEXT: // %bb.0:64; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cluster_space_cta_param_0];65; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_arrive_scope_cluster_space_cta_param_1];66; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.release.cluster.shared.b64 %rd1, [%r1], %r2;67; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.expect_tx.release.cluster.shared.b64 %rd2, [%r1], %r2;68; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.release.cluster.shared.b64 %rd3, [%r1], %r2;69; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.expect_tx.release.cluster.shared.b64 %rd4, [%r1], %r2;70; CHECK-PTX-SHARED32-NEXT: ret;71 %r1 = call i64 @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx)72 %r2 = call i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx)73 %r3 = call i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx)74 %r4 = call i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %tx)75 ret void76}77 78; === space_cluster (addrspace 7) ===79define void @test_mbarrier_arrive_scope_cta_space_cluster(ptr addrspace(7) %mbar, i32 %tx) {80; CHECK-PTX64-LABEL: test_mbarrier_arrive_scope_cta_space_cluster(81; CHECK-PTX64: {82; CHECK-PTX64-NEXT: .reg .b32 %r<2>;83; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;84; CHECK-PTX64-EMPTY:85; CHECK-PTX64-NEXT: // %bb.0:86; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_mbarrier_arrive_scope_cta_space_cluster_param_0];87; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cta_space_cluster_param_1];88; CHECK-PTX64-NEXT: mbarrier.arrive.shared::cluster.b64 _, [%rd1], %r1;89; CHECK-PTX64-NEXT: mbarrier.arrive.expect_tx.shared::cluster.b64 _, [%rd1], %r1;90; CHECK-PTX64-NEXT: mbarrier.arrive_drop.shared::cluster.b64 _, [%rd1], %r1;91; CHECK-PTX64-NEXT: mbarrier.arrive_drop.expect_tx.shared::cluster.b64 _, [%rd1], %r1;92; CHECK-PTX64-NEXT: ret;93;94; CHECK-PTX-SHARED32-LABEL: test_mbarrier_arrive_scope_cta_space_cluster(95; CHECK-PTX-SHARED32: {96; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;97; CHECK-PTX-SHARED32-EMPTY:98; CHECK-PTX-SHARED32-NEXT: // %bb.0:99; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cta_space_cluster_param_0];100; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_arrive_scope_cta_space_cluster_param_1];101; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.shared::cluster.b64 _, [%r1], %r2;102; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.expect_tx.shared::cluster.b64 _, [%r1], %r2;103; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.shared::cluster.b64 _, [%r1], %r2;104; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.expect_tx.shared::cluster.b64 _, [%r1], %r2;105; CHECK-PTX-SHARED32-NEXT: ret;106 call void @llvm.nvvm.mbarrier.arrive.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx)107 call void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx)108 call void @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx)109 call void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %mbar, i32 %tx)110 ret void111}112 113define void @test_mbarrier_arrive_scope_cluster_space_cluster(ptr addrspace(7) %mbar, i32 %tx) {114; CHECK-PTX64-LABEL: test_mbarrier_arrive_scope_cluster_space_cluster(115; CHECK-PTX64: {116; CHECK-PTX64-NEXT: .reg .b32 %r<2>;117; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;118; CHECK-PTX64-EMPTY:119; CHECK-PTX64-NEXT: // %bb.0:120; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [test_mbarrier_arrive_scope_cluster_space_cluster_param_0];121; CHECK-PTX64-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cluster_space_cluster_param_1];122; CHECK-PTX64-NEXT: mbarrier.arrive.release.cluster.shared::cluster.b64 _, [%rd1], %r1;123; CHECK-PTX64-NEXT: mbarrier.arrive.expect_tx.release.cluster.shared::cluster.b64 _, [%rd1], %r1;124; CHECK-PTX64-NEXT: mbarrier.arrive_drop.release.cluster.shared::cluster.b64 _, [%rd1], %r1;125; CHECK-PTX64-NEXT: mbarrier.arrive_drop.expect_tx.release.cluster.shared::cluster.b64 _, [%rd1], %r1;126; CHECK-PTX64-NEXT: ret;127;128; CHECK-PTX-SHARED32-LABEL: test_mbarrier_arrive_scope_cluster_space_cluster(129; CHECK-PTX-SHARED32: {130; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;131; CHECK-PTX-SHARED32-EMPTY:132; CHECK-PTX-SHARED32-NEXT: // %bb.0:133; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [test_mbarrier_arrive_scope_cluster_space_cluster_param_0];134; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [test_mbarrier_arrive_scope_cluster_space_cluster_param_1];135; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.release.cluster.shared::cluster.b64 _, [%r1], %r2;136; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive.expect_tx.release.cluster.shared::cluster.b64 _, [%r1], %r2;137; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.release.cluster.shared::cluster.b64 _, [%r1], %r2;138; CHECK-PTX-SHARED32-NEXT: mbarrier.arrive_drop.expect_tx.release.cluster.shared::cluster.b64 _, [%r1], %r2;139; CHECK-PTX-SHARED32-NEXT: ret;140 call void @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx)141 call void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx)142 call void @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx)143 call void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %mbar, i32 %tx)144 ret void145}146 147declare i64 @llvm.nvvm.mbarrier.arrive.scope.cta.space.cta(ptr addrspace(3), i32)148declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cta(ptr addrspace(3), i32)149declare i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta(ptr addrspace(3), i32)150declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cta(ptr addrspace(3), i32)151 152declare i64 @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cta(ptr addrspace(3), i32)153declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cta(ptr addrspace(3), i32)154declare i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cta(ptr addrspace(3), i32)155declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cta(ptr addrspace(3), i32)156 157declare void @llvm.nvvm.mbarrier.arrive.scope.cta.space.cluster(ptr addrspace(7), i32)158declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cluster(ptr addrspace(7), i32)159declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cluster(ptr addrspace(7), i32)160declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cluster(ptr addrspace(7), i32)161 162declare void @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cluster(ptr addrspace(7), i32)163declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cluster(ptr addrspace(7), i32)164declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cluster(ptr addrspace(7), i32)165declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cluster(ptr addrspace(7), i32)166