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