diff options
| author | Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com> | 2016-11-15 19:00:15 +0000 |
|---|---|---|
| committer | Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com> | 2016-11-15 19:00:15 +0000 |
| commit | ea91cca593bb543edce86c15c2948ef71b022932 (patch) | |
| tree | aea58698cfc016852db61f0eedd5f8ffa1a4353a /llvm/include | |
| parent | cd433d28118b345ba367420cbed1acd8fb1e4ea3 (diff) | |
| download | bcm5719-llvm-ea91cca593bb543edce86c15c2948ef71b022932.tar.gz bcm5719-llvm-ea91cca593bb543edce86c15c2948ef71b022932.zip | |
[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
Diffstat (limited to 'llvm/include')
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 3 |
1 files changed, 3 insertions, 0 deletions
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< |
