llvm/test/CodeGen/NVPTX/barrier.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | %ptxas-verify %}

declare void @llvm.nvvm.bar.warp.sync(i32)
declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32)
declare void @llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32)
declare void @llvm.nvvm.barrier.cta.sync.all(i32)
declare void @llvm.nvvm.barrier.cta.sync.count(i32, i32)
declare void @llvm.nvvm.barrier.cta.arrive.count(i32, i32)
declare void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32)

define void @barrier_warp_sync(i32 %id) {
; CHECK-LABEL: barrier_warp_sync(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_warp_sync_param_0];
; CHECK-NEXT:    bar.warp.sync %r1;
; CHECK-NEXT:    bar.warp.sync 6;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.bar.warp.sync(i32 %id)
  call void @llvm.nvvm.bar.warp.sync(i32 6)
  ret void
}

define void @barrier_cta_sync_aligned_all(i32 %id) {
; CHECK-LABEL: barrier_cta_sync_aligned_all(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_sync_aligned_all_param_0];
; CHECK-NEXT:    bar.sync %r1;
; CHECK-NEXT:    bar.sync 3;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id)
  call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 3)
  ret void
}

define void @barrier_cta_sync_aligned(i32 %id, i32 %cnt) {
; CHECK-LABEL: barrier_cta_sync_aligned(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_sync_aligned_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [barrier_cta_sync_aligned_param_1];
; CHECK-NEXT:    bar.sync %r1, %r2;
; CHECK-NEXT:    bar.sync 3, %r2;
; CHECK-NEXT:    bar.sync %r1, 64;
; CHECK-NEXT:    bar.sync 4, 64;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 3, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 64)
  call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 4, i32 64)
  ret void
}

define void @barrier_cta_arrive_aligned(i32 %id, i32 %cnt) {
; CHECK-LABEL: barrier_cta_arrive_aligned(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_arrive_aligned_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [barrier_cta_arrive_aligned_param_1];
; CHECK-NEXT:    bar.arrive %r1, %r2;
; CHECK-NEXT:    bar.arrive 3, %r2;
; CHECK-NEXT:    bar.arrive %r1, 64;
; CHECK-NEXT:    bar.arrive 4, 64;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 3, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 64)
  call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 4, i32 64)
  ret void
}

define void @barrier_cta_sync_all(i32 %id) {
; CHECK-LABEL: barrier_cta_sync_all(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_sync_all_param_0];
; CHECK-NEXT:    barrier.sync %r1;
; CHECK-NEXT:    barrier.sync 3;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.sync.all(i32 %id)
  call void @llvm.nvvm.barrier.cta.sync.all(i32 3)
  ret void
}

define void @barrier_cta_sync(i32 %id, i32 %cnt) {
; CHECK-LABEL: barrier_cta_sync(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_sync_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [barrier_cta_sync_param_1];
; CHECK-NEXT:    barrier.sync %r1, %r2;
; CHECK-NEXT:    barrier.sync 3, %r2;
; CHECK-NEXT:    barrier.sync %r1, 64;
; CHECK-NEXT:    barrier.sync 4, 64;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.sync.count(i32 3, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 64)
  call void @llvm.nvvm.barrier.cta.sync.count(i32 4, i32 64)
  ret void
}

define void @barrier_cta_arrive(i32 %id, i32 %cnt) {
; CHECK-LABEL: barrier_cta_arrive(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [barrier_cta_arrive_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [barrier_cta_arrive_param_1];
; CHECK-NEXT:    barrier.arrive %r1, %r2;
; CHECK-NEXT:    barrier.arrive 3, %r2;
; CHECK-NEXT:    barrier.arrive %r1, 64;
; CHECK-NEXT:    barrier.arrive 4, 64;
; CHECK-NEXT:    ret;
  call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.arrive.count(i32 3, i32 %cnt)
  call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 64)
  call void @llvm.nvvm.barrier.cta.arrive.count(i32 4, i32 64)
  ret void
}