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) 41
