137 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s3; RUN: %if ptxas-isa-6.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | %ptxas-verify %}4 5declare void @llvm.nvvm.bar.warp.sync(i32)6declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32)7declare void @llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32)8declare void @llvm.nvvm.barrier.cta.sync.all(i32)9declare void @llvm.nvvm.barrier.cta.sync.count(i32, i32)10declare void @llvm.nvvm.barrier.cta.arrive.count(i32, i32)11declare void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32)12 13define void @barrier_warp_sync(i32 %id) {14; CHECK-LABEL: barrier_warp_sync(15; CHECK: {16; CHECK-NEXT: .reg .b32 %r<2>;17; CHECK-EMPTY:18; CHECK-NEXT: // %bb.0:19; CHECK-NEXT: ld.param.b32 %r1, [barrier_warp_sync_param_0];20; CHECK-NEXT: bar.warp.sync %r1;21; CHECK-NEXT: bar.warp.sync 6;22; CHECK-NEXT: ret;23 call void @llvm.nvvm.bar.warp.sync(i32 %id)24 call void @llvm.nvvm.bar.warp.sync(i32 6)25 ret void26}27 28define void @barrier_cta_sync_aligned_all(i32 %id) {29; CHECK-LABEL: barrier_cta_sync_aligned_all(30; CHECK: {31; CHECK-NEXT: .reg .b32 %r<2>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_sync_aligned_all_param_0];35; CHECK-NEXT: bar.sync %r1;36; CHECK-NEXT: bar.sync 3;37; CHECK-NEXT: ret;38 call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id)39 call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 3)40 ret void41}42 43define void @barrier_cta_sync_aligned(i32 %id, i32 %cnt) {44; CHECK-LABEL: barrier_cta_sync_aligned(45; CHECK: {46; CHECK-NEXT: .reg .b32 %r<3>;47; CHECK-EMPTY:48; CHECK-NEXT: // %bb.0:49; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_sync_aligned_param_0];50; CHECK-NEXT: ld.param.b32 %r2, [barrier_cta_sync_aligned_param_1];51; CHECK-NEXT: bar.sync %r1, %r2;52; CHECK-NEXT: bar.sync 3, %r2;53; CHECK-NEXT: bar.sync %r1, 64;54; CHECK-NEXT: bar.sync 4, 64;55; CHECK-NEXT: ret;56 call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 %cnt)57 call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 3, i32 %cnt)58 call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 64)59 call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 4, i32 64)60 ret void61}62 63define void @barrier_cta_arrive_aligned(i32 %id, i32 %cnt) {64; CHECK-LABEL: barrier_cta_arrive_aligned(65; CHECK: {66; CHECK-NEXT: .reg .b32 %r<3>;67; CHECK-EMPTY:68; CHECK-NEXT: // %bb.0:69; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_arrive_aligned_param_0];70; CHECK-NEXT: ld.param.b32 %r2, [barrier_cta_arrive_aligned_param_1];71; CHECK-NEXT: bar.arrive %r1, %r2;72; CHECK-NEXT: bar.arrive 3, %r2;73; CHECK-NEXT: bar.arrive %r1, 64;74; CHECK-NEXT: bar.arrive 4, 64;75; CHECK-NEXT: ret;76 call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 %cnt)77 call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 3, i32 %cnt)78 call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 64)79 call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 4, i32 64)80 ret void81}82 83define void @barrier_cta_sync_all(i32 %id) {84; CHECK-LABEL: barrier_cta_sync_all(85; CHECK: {86; CHECK-NEXT: .reg .b32 %r<2>;87; CHECK-EMPTY:88; CHECK-NEXT: // %bb.0:89; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_sync_all_param_0];90; CHECK-NEXT: barrier.sync %r1;91; CHECK-NEXT: barrier.sync 3;92; CHECK-NEXT: ret;93 call void @llvm.nvvm.barrier.cta.sync.all(i32 %id)94 call void @llvm.nvvm.barrier.cta.sync.all(i32 3)95 ret void96}97 98define void @barrier_cta_sync(i32 %id, i32 %cnt) {99; CHECK-LABEL: barrier_cta_sync(100; CHECK: {101; CHECK-NEXT: .reg .b32 %r<3>;102; CHECK-EMPTY:103; CHECK-NEXT: // %bb.0:104; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_sync_param_0];105; CHECK-NEXT: ld.param.b32 %r2, [barrier_cta_sync_param_1];106; CHECK-NEXT: barrier.sync %r1, %r2;107; CHECK-NEXT: barrier.sync 3, %r2;108; CHECK-NEXT: barrier.sync %r1, 64;109; CHECK-NEXT: barrier.sync 4, 64;110; CHECK-NEXT: ret;111 call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 %cnt)112 call void @llvm.nvvm.barrier.cta.sync.count(i32 3, i32 %cnt)113 call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 64)114 call void @llvm.nvvm.barrier.cta.sync.count(i32 4, i32 64)115 ret void116}117 118define void @barrier_cta_arrive(i32 %id, i32 %cnt) {119; CHECK-LABEL: barrier_cta_arrive(120; CHECK: {121; CHECK-NEXT: .reg .b32 %r<3>;122; CHECK-EMPTY:123; CHECK-NEXT: // %bb.0:124; CHECK-NEXT: ld.param.b32 %r1, [barrier_cta_arrive_param_0];125; CHECK-NEXT: ld.param.b32 %r2, [barrier_cta_arrive_param_1];126; CHECK-NEXT: barrier.arrive %r1, %r2;127; CHECK-NEXT: barrier.arrive 3, %r2;128; CHECK-NEXT: barrier.arrive %r1, 64;129; CHECK-NEXT: barrier.arrive 4, 64;130; CHECK-NEXT: ret;131 call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 %cnt)132 call void @llvm.nvvm.barrier.cta.arrive.count(i32 3, i32 %cnt)133 call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 64)134 call void @llvm.nvvm.barrier.cta.arrive.count(i32 4, i32 64)135 ret void136}137