test/CodeGen/NVPTX/named-barriers.ll

   1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
   2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
   3
   4 ; Use bar.sync to arrive at a pre-computed barrier number and
   5 ; wait for all threads in CTA to also arrive:
   6 define ptx_device void @test_barrier_named_cta() {
   7 ; CHECK: mov.u32  %r[[REG0:[0-9]+]], 0;
   8 ; CHECK: bar.sync %r[[REG0]];
   9 ; CHECK: mov.u32  %r[[REG1:[0-9]+]], 10;
  10 ; CHECK: bar.sync %r[[REG1]];
  11 ; CHECK: mov.u32  %r[[REG2:[0-9]+]], 15;
  12 ; CHECK: bar.sync %r[[REG2]];
  13 ; CHECK: ret;
  14   call void @llvm.nvvm.barrier.n(i32 0)
  15   call void @llvm.nvvm.barrier.n(i32 10)
  16   call void @llvm.nvvm.barrier.n(i32 15)
  17   ret void
  18 }
  19
  20 ; Use bar.sync to arrive at a pre-computed barrier number and
  21 ; wait for fixed number of cooperating threads to arrive:
  22 define ptx_device void @test_barrier_named() {
  23 ; CHECK: mov.u32  %r[[REG0A:[0-9]+]], 32;
  24 ; CHECK: mov.u32  %r[[REG0B:[0-9]+]], 0;
  25 ; CHECK: bar.sync %r[[REG0B]], %r[[REG0A]];
  26 ; CHECK: mov.u32  %r[[REG1A:[0-9]+]], 352;
  27 ; CHECK: mov.u32  %r[[REG1B:[0-9]+]], 10;
  28 ; CHECK: bar.sync %r[[REG1B]], %r[[REG1A]];
  29 ; CHECK: mov.u32  %r[[REG2A:[0-9]+]], 992;
  30 ; CHECK: mov.u32  %r[[REG2B:[0-9]+]], 15;
  31 ; CHECK: bar.sync %r[[REG2B]], %r[[REG2A]];
  32 ; CHECK: ret;
  33   call void @llvm.nvvm.barrier(i32 0, i32 32)
  34   call void @llvm.nvvm.barrier(i32 10, i32 352)
  35   call void @llvm.nvvm.barrier(i32 15, i32 992)
  36   ret void
  37 }
  38
  39 declare void @llvm.nvvm.barrier(i32, i32)
  40 declare void @llvm.nvvm.barrier.n(i32)