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:
6define 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:
22define 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
39declare void @llvm.nvvm.barrier(i32, i32)
40declare void @llvm.nvvm.barrier.n(i32)
41