From ea91cca593bb543edce86c15c2948ef71b022932 Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Tue, 15 Nov 2016 19:00:15 +0000 Subject: [AMDGPU] Add wave barrier builtin The wave barrier represents the discardable barrier. Its main purpose is to carry convergent attribute, thus preventing illegal CFG optimizations. All lanes in a wave come to convergence point simultaneously with SIMT, thus no special instruction is needed in the ISA. The barrier is discarded during code generation. Differential Revision: https://reviews.llvm.org/D26585 llvm-svn: 287007 --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 3 +++ 1 file changed, 3 insertions(+) (limited to 'llvm/include') diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index cc4fd4ce5f4..5105e0d68ae 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -107,6 +107,9 @@ def int_amdgcn_dispatch_id : def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">, Intrinsic<[], [], [IntrConvergent]>; +def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">, + Intrinsic<[], [], [IntrConvergent]>; + def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>; def int_amdgcn_div_scale : Intrinsic< -- cgit v1.2.3